CUDA估計每塊和塊編號線程2D網格數據

Question

讓我說，我讀過關於SO仔細閱讀所有類似的問題開始：

1. Determining threads per block and block per grid
2. Threads per SM, threads per block
3. CUDA Blocks and Threads
4. Warps and optimal number of blocks

我的意圖是嘗試和計算dynami爲我正在開發的前饋神經網絡庫提供（而不是硬編碼值）。

我的數據不是正方形格子（矩陣狀），這是常有與我見過最實施例中，它是不是兩個矢量產生的矩陣，具有不相等的行列：

float x[6] {1.f, 1.f, 0.f, 1.f, 1.f, 0.f}; 
thrust::device_vector<float> in_vec(x, x+6); 
float y[9] {1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f}; 
thrust::device_vector<float> w_vec(y, y+9); 
thrust::device_vector<float> o_wec(9); 
thrust::device_vector<float> mtx_vec(9 * 6); 

float * i_ptr = thrust::raw_pointer_cast(in_vec.data()); 
float * w_ptr = thrust::raw_pointer_cast(w_vec.data()); 
float * out_ptr = thrust::raw_pointer_cast(mtx_vec.data()); 

dim3 threadsPerBlock(9,6); 
dim3 numBlocks(1,1); 
prop_mtx<<<numBlocks,threadsPerBlock>>>(w_ptr, i_ptr, out_ptr, 6); 

和內核：

__global__ void prop_mtx(float * w, float * i, float * o, int s) 
{ 
    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 
    o[y + x * s] = w[x] * i[y]; 
} 

爲什麼我已經採取了這種做法，因爲它使神經網絡計算的意義，當它涉及到向量/矩陣計算的原因。 我想保持這個一致性，並且使用用於權重*輸入計算的2D網格的AFAIK是合理的。

我必須將每個塊的線程計算爲2D，並且網格中的線程數量不等。

我尤斯一個GTX 660，其中有：

CUDA Capability Major/Minor version number: 3.0 
    Total amount of global memory:     2047 MBytes 
    (5) Multiprocessors, (192) CUDA Cores/MP:  960 CUDA Cores 
    Maximum Texture Dimension Size (x,y,z)   1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096) 
    Warp size:          32 
    Maximum number of threads per multiprocessor: 2048 
    Maximum number of threads per block:   1024 
    Max dimension size of a thread block (x,y,z): (1024, 1024, 64) 
    Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) 

我想明白，我怎麼可以推斷/計算網格尺寸，每塊的線程和塊的數量。

讓我們假設我有一個800個項目的權重向量和一個6500個項目的輸入向量。

1. 這是否意味着我真正需要的是一個800,6500的二維網格？據我所知，其他的都會提供不正確的結果？

我知道我的最大線程每塊爲1024，但由於其在2D網格，它更可能是：

dim3 threadPerBlock(X,Y); 

由於事實上，我的網格不是一個方形矩陣，我需要以不同的方式計算每個塊的X，Y線程？

或者我需要先推斷出所需的塊數？

最後，由於我的螺紋經大小爲32，

是否最小網格大小，而不管所有其它參數需要爲至少32 ，還是32的倍數？我是否需要至少每塊32個線程，或最小數目爲32的網格大小？

任何僞代碼，或解釋我應該如何去做，將不勝感激。

我試過的是計算我的2D網格大小，通過將我的數據除以32包裝大小。 然後我考慮使用可用的SM計算網格線程。例如

800 weights/5 SM, = 160 x's per SM 
6500 inputs/5 SM, = 1300 y's per SM 

但我不知道該從那裏做什麼。 最後，我認爲首先找到輸入重量比：

6500/800 = 8.125 

言下之意是使用X的32最小網格大小， ÿ將不得不由8.125 * 32 因此相乘，我的threadsPerBlock將是：

dim3 threadsPerBlock(32,260); 

那當然，每塊有8320個線程，遠遠超過每塊1024個。

所以這是我的問題：我怎麼不超過每塊1024線程，同時保留我的數據正確的網格大小？

PS：我的問題不是優化代碼，而是瞭解如何在設備上分配線程和網格數據。

Answer 1

分類計算問題的一種方法是討論轉換和減少。

A 減少是一類問題，它需要一個大的輸入數據集大小，併產生一個小的輸出數據集大小。例如，拍攝圖像並找到最大像素值將是一個減少。對於這個討論，我們將忽略減少。

A 轉換是輸出數據集大小（元素數）與輸入數據集大小「大」或「大致相同」的計算類別。例如，拍攝圖像並生成模糊圖像將是一種轉變。

對於轉換，編寫cuda內核（線程代碼）的常用方法（「線程策略」）將是爲輸出數組中的每個點創建一個唯一的線程。因此，我必須擁有的最小線程總數等於我輸出數組的大小。線程代碼只是輸入數據所需的一組計算，以便產生一個輸出數據點。粗略地說，那麼你的問題和簡化的內核就符合這個定義;這是一個轉變。

遵循上述線程策略，我們需要網格中的線程總數等於我需要創建的輸出點總數。對於二維問題，考慮這些二維問題通常比較方便，CUDA爲此提供二維（或三維）線程塊組織和二維（或三維）網格組織。

CUDA線程塊維度的選擇通常有些隨意。一般來說，我們通常希望針對每個塊範圍內128到512個線程的線程塊（出於其他地方所述的原因），並且我們希望線程塊獲得32的整數倍數（warp大小）的線程塊以提高效率細分爲經線，這是CUDA執行的實際單位。在當前支持的GPU上，線程塊被限制爲每塊1024線程（總數 - 即尺寸的乘積）。然而，對於許多問題，這個範圍內的線程塊選擇（例如，256個線程與512個線程）通常對性能影響相對較小。爲了獲得某種工作的興趣，我們現在不會在這方面出現細節。 （當你回來進行優化時，你可能會重新考慮這個選擇。）

到目前爲止，我們已經知道對於這個問題類型，我們需要線程的總數來覆蓋我們的問題空間，有點任意的線程塊尺寸選擇。所以讓我們選擇（32,16）（x，y）開始，總共512個線程。沒有任何規則指出，新區塊需要是「方形」，或者網格需要是「方形」，或者甚至應該存在線程塊尺寸和問題尺寸（或網格尺寸）之間的任何比率校驗。

現在我們有了（32,16）的選擇，我們必須問自己「我需要多少這些？」。這個問題是2D的，所以我們選擇了一個2D線程塊來簡化線程代碼中的索引生成。我們還要選擇一個2D網格 - 它對2D問題很有意義，而且對於索引生成的2D簡單性來說也是如此。所以我們可以獨立考慮這兩個維度。

那麼，我需要在x方向上需要多少塊？我至少需要（x中的問題大小）/（我在x中的線程塊大小）。既然我們在這裏處理所有的整數，這就引出了一個問題：「如果我的問題大小不能被我的線程塊大小整除？」規範的解決方案是發射綽綽有餘的線程來覆蓋空間，或者足夠的線程來覆蓋空間。但在非均勻分佈的情況下，這將導致「額外的線程」。我們將盡快討論並處理這些問題。因此，如果我有threadblock尺寸這樣一個爲dim3變量：

#define BX 32 
    #define BY 16 
    ... 
    dim3 block(BX,BY); 

那麼我會建立我爲dim3格變量是這樣的：如果你通過這個運算工作

#define DX 800 
    #define DY 6500 
    ... 
    dim3 grid((DX+block.x-1)/block.x, (DY+block.y-1)/block.y); 

，你會看到，這導致我們在x和y方向上啓動足夠的塊，以便我們至少有足夠的線程來覆蓋（DX，DY）的問題空間，每個輸出點一個線程。

希望很明顯，Y維度是分開處理和獨立於x維度的。

上面的計算通常會導致在我的網格中生成「太多」線程。除了我需要處理的問題空間（DX，DY）之外，我還會有一些「額外的線程」。我們希望這些線程「無所事事」。處理這個問題的規範方法是將問題空間維度傳遞給我的內核，在我的內核中創建適當的全局唯一線程索引，然後將該索引與我的問題空間中的最大索引進行比較。如果它超過它，我們只需要該線程跳過所有剩餘的線程代碼。

使用內核爲例，它可能是這樣的：即在隨後的代碼「不參與」

__global__ void prop_mtx(float * w, float * i, float * o, int s, const size_t d_size_x, const size_t d_size_y) 
{ 
    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 
    if ((x < d_size_x) && (y < d_size_y)) // thread check 
     o[y + x * s] = w[x] * i[y]; 
} 

注意這樣一個線程檢查將創建線程（在某些塊）。需要注意的一點是，__syncthreads()的使用取決於參與塊中的所有線程。因此，在這種情況下，我們不應該直接使用__syncthreads()。取而代之的是，我們要適當地調節threadblock行爲：

__global__ void prop_mtx(float * w, float * i, float * o, int s, const size_t d_size_x, const size_t d_size_y) 
{ 
    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 
    if ((x < d_size_x) && (y < d_size_y)) // thread check 
     { 
     o[y + x * s] = w[x] * i[y]; 
     // and other code not dependent on __syncthreads() 
     } 
    // now it is safe to use since all threads are participating 
    __syncthreads(); 
    if ((x < d_size_x) && (y < d_size_y)) // thread check 
     { 
      // rest of kernel code 
     } 
} 

注意，有可能具有的線程數量較少進行必要的計算輸出數據點的數量較多。線程和輸出數據之間的1：1對應關係是思考和編寫cuda內核代碼的簡單方法，但這不是唯一的方法。另外一種可能的方法是使用某種形式的網格跨步迴路，這樣較小的網格可以覆蓋更大的問題空間。對這些策略的討論超出了這個答案的範圍，應該在理解其他方法之前理解在這個答案中討論的基本方法。