有没有更好的方法来处理 CUDA 中的"undividable count of numbers by block_size"？

Question

我需要对N个数的向量进行数据约简（找到k-max数）。问题是，我事先（在编译之前）不知道N，当我构建两个内核时，我也不确定我做得是否正确——一个内核有(int)(N / block_size)块，第二个内核有一个N % block_size线程块。

是否有更好的方法来处理CUDA中按block_size"不可分割"的数字计数

Answer 1

一个典型的方法如下（1-D网格示例）：

#define DATA_SIZE ...   // this is some arbitrary number
#define NUM_THREADS_PER_BLOCK ...  // this is block size, usually a multiple of 32
                                  // typical choices are 256, 512, 1024 (pick one)
unsigned int N = DATA_SIZE;  
unsigned int nTPB = NUM_THREADS_PER_BLOCK; 
my_kernel<<<(N + nTPB - 1)/nTPB, nTPB>>>(...);

这假设您的内核在开始时有一个"线程检查"，如下所示：

unsigned int idx = threadIdx.x + blockDim.x*blockIdx.x;
if (idx < DATA_SIZE){
   // kernel code goes here
}

Answer 2

@RobertCrovella的回答描述了处理这种情况的标准方法，通常不需要担心内核中需要的额外if条件。

然而，另一种选择是为输入和输出缓冲区分配最多可被块大小整除的填充数，运行内核（不带if），然后忽略额外的结果，例如不将它们复制回CPU。