從給定的二維偏移量知道CUDA中的塊ID

Question

我試圖從給定的CUDA偏移量計算blockIdx.x和blockIdx.y，但我完全被大腦阻塞。這個想法是在可能的情況下從共享內存讀取數據，在其他情況下是從全局內存中讀取數據

在例如，如果我已經64個元件一維陣列和我配置具有16X1螺紋的內核（總共4個塊），每個線程可以使用訪問的位置：

int idx = blockDim.x*blockIdx.x + threadIdx.x 

，我可以容易從IDX得到一個給定的索引值的blockIdx.x作爲

int blockNumber = idx/blockDim.x; 

但與8×8元素和4×4的線程的內核配置（2×2塊總共）二維場景每個線程使用訪問的位置：

int x = threadIdx.x + blockIdx.x * blockDim.x; 
int y = threadIdx.y + blockIdx.y * blockDim.y; 
int pitch = blockDim.x * gridDim.x; 
int idx = x + y * pitch; 

int sharedMemIndex = threadIdx.x+threadIdx.y+BLOCK_DIM_X; 
__shared_block[sharedMemIndex] = fromGlobalMemory[idx]; 
__syncthreads(); 

// ... some operations 

int unknow_index = __shared_block[sharedMemIndex]; 

if (unknow_index within this block?) 
    // ... read from shared memory 
else 
    // ... read from global memory 

我該如何知道給定idx上的Block ID.x和ID.y？即索引34和35在塊（1,1）中的塊（0,1）和索引36中。因此，如果塊（0,1）中的線程讀取索引35的值，該線程將知道該值位於其塊中，並將從共享內存中讀取該值。索引35值將被存儲在塊的共享存儲器的位置11（0.1）中。

在此先感謝！

Answer 1

在實踐中，我實在想不出一個很好的理由，這是以往任何時候都需要的，但你可以計算出這樣的結果，對於任意索引值idx（假設列排序索引）：

int pitch = blockDim.x * gridDim.x; 
int tidy = idx/pitch; // div(idx,pitch) 
int tidx = idx - (pitch * tidy); // mod(idx,pitch) 
int bidx = idx/blockDim.x; 
int bidy = idy/blockDim.y; 

它應該爲您提供bidx和bidy中索引的塊座標。

Answer 2

您正在執行不必要的計算。

idx/blockDim.x 
-->(blockDim.x * blockIdx.x + threadIdx.x)/blockDim.x 
-->(blockIdx.x + threadIdx.x/blockDim.x) 
--> blockIdx.x + 0 (threadIdx.x always less than blockDim.x) 

您可以使用blockIdx.x代替卷積計算。 2D網格也是如此（blockIdx.x和blockIdx.y）。

Answer 3

有沒有必要應用Idx的數學來找出X和Y塊或從Idx向後查找塊索引。對於每個線程（Idx），只需調用blockIdx.x和blockIdx.y即可找出Y和X塊。

在內核中的任何一點：

int x = blockIdx.x // will give you X block Index at that particular thread 
int y = blockIdx.y // will give you Y block Index at that particular thread. 

更新： 如果你在反向操作死心塌地，你需要知道的間距和塊尺寸的值

int currentRow = idx/pitch; 
    int currentCol = idx%pitch; 

    int block_idx_x = currentCol/blockDim.x; 
    int block_idx_y = currentRow/blockDim.y;