CUDA求和归约难题

Question

CUDA的缩减完全把我搞糊涂了！首先，Mark Harris的this教程和Mike Giles的this one都使用了声明extern __shared__ temp[]。关键字extern在C中用于声明，但分配发生在"elsewhre“(例如，通常在另一个C文件上下文中)。这里extern的相关性是什么？为什么我们不使用：

__shared__ float temp[N/2];

例如?或者我们为什么不将temp声明为全局变量，例如

#define N 1024
__shared__ float temp[N/2];

__global__ void sum(float *sum,  float *data){ ... }

int main(){
 ...
 sum<<<M,L>>>(sum, data);
}

我还有另一个问题？应该使用多少个块和每个块的线程来调用求和内核？我尝试了this example (基于this)。

注:您可以找到有关my devices here的信息。

Answer 1

第一个问题的答案是，CUDA支持在运行时动态分配共享内存(有关更多详细信息，请参阅this SO memory和the documentation )。使用extern声明共享内存向编译器表明，共享内存大小将在内核启动时确定，并以字节作为参数传递给<<< >>>语法(或等效地通过API函数)，类似于：

sum<<< gridsize, blocksize, sharedmem_size >>>(....);

第二个问题通常是启动块的数量，这将完全填满您的GPU上的所有流多处理器。大多数合理编写的缩减内核将在每个线程中累积许多值，然后执行共享内存缩减。减少要求每个块的线程数是2的幂:这通常是32，64,128,256,512 (或者1024，如果你有费米或开普勒GPU)。这是一个非常有限的搜索空间，只是作为基准来看看什么在你的硬件上工作得最好。您可以找到关于块和网格大小调整here和here的更一般的讨论。