从给定的二维偏移量知道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;