理解CUDA核函数的启动参数

Question

我正在尝试分析一些我在网上找到的代码，但我一直想不通。我正在查看使用以下参数启动的直方图内核

histogram<<<2500, numBins, numBins * sizeof(unsigned int)>>>(...); 

我知道参数是网格、块、共享内存大小。

那是否意味着有2500个块，每个块都有numBins个线程，每个块还有一个numBins * sizeof(unsigned int)大小的共享内存块可用于其线程？

此外，在内核内部也会调用__syncthreads()，那么在内核调用过程中会有2500组numBins个对__syncthreads()的调用吗？

- KDecker

---

第一个问题：是的。 第二个问题：在那2500个块内部的线程，与其他块中的线程无关，会到达__syncthreads()点，等待直到该块内所有线程更新共享内存到该点并到达，然后继续执行后续指令。该块中的所有线程都必须看到__syncthreads()，因此可以说有2500组对numBins的__syncthreads()调用，但这与普通函数调用不同。 这是用于块内线程同步的屏障例程。 - Farzad

Answer 1

那是否意味着每个numBins线程块都有2500个块，每个块还有一个大小为numBins * sizeof(unsigned int)的共享内存块可供其线程使用？

来自CUDA Toolkit documentation：

全局函数调用的执行配置由插入以下形式的表达式指定：<<>>,其中：

- Dg（dim3）指定网格的维度和大小。 - Db（dim3）指定每个块的维度和大小 - Ns（size_t）指定在此调用中为每个块动态分配的共享内存中的字节数，除了静态分配的内存之外。 - S（cudaStream_t）指定关联的流，是可选参数，默认为0。

所以，正如@Fazar指出的那样，答案是肯定的。这个内存是按块分配的。

此外，在内核本身中有对__syncthreads()的调用，那么在内核调用过程中会有2500组numBins对__syncthreads()的调用吗？

__syncthreads()等待所有线程到达此点。用于协调同一块中线程之间的通信。

因此，每个块都有一个__syncthread()调用。

- srodrb