CUDA 求和归约之谜

Question

CUDA 的减少让我完全感到困惑！首先，Mark Harris 的 this 教程和 Mike Giles 的 this one 都使用了声明 extern __shared__ temp[]。声明时在 C 中使用关键字 extern，但分配发生在“elsehre”（例如，通常在另一个 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）。

注意：您可以找到有关我的设备的信息here。

Answer 1

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

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

第二个问题通常是启动将完全填满 GPU 上所有流式多处理器的块数。大多数明智地编写的缩减内核将在每个线程中累积许多值，然后执行共享内存缩减。减少要求每个块的线程数是 2 的幂：这通常为您提供 32、64、128、256、512（如果您有 Fermi 或 Kepler GPU，则为 1024）。这是一个非常有限的搜索空间，只需进行基准测试，看看什么在您的硬件上最有效。您可以在here 和here 找到有关块和网格大小的更一般性的讨论。