GPGPU：一维线程索引的 CUDA 内核配置 - 线程、块、共享内存和寄存器

Question

假设我有 N 个任务，其中每个任务都可以由 GPU 上的单个线程执行。还假设 N = GPU 上的线程数。

问题一： 以下是启动 maximum 大小的一维内核的合适方法吗？ GPU 上存在的所有 N 个线程都会执行这项工作吗？

cudaDeviceProp  theProps;

dim3 mygrid(theProps.maxGridSize[0], 1, 1);
dim3 myblocks(theProps.maxThreadsDim[0], 1, 1);

mykernel<<<mygrid, myblocks>>>(...);

问题2：

cudaDeviceProp::maxThreadsPerBlock

与

cudaDeviceProp::maxThreadsDim[0]

的关系是什么？它们有何不同？

cudaDeviceProp::maxThreadsPerBlock

可以代替上面的

cudaDeviceProp::maxThreadsDim[0]

吗？

问题3： 假设我想在块中的线程之间平均分配块的共享内存，并且我希望每个线程都可以使用最大数量的共享内存。那我应该最大化块的数量，并最小化每个块的线程数，对吧？

问题4： 只是为了确认（在查看有关 SO 的相关问题之后），在上面的线性（1D）网格/块方案中，全局唯一线程索引是

unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x

。对吧？

Answer 1

建议每个问题问一个问题。面对各种各样的问题，任何人都很难给出完整的答案。 SO 并不是真正的教程服务。您应该利用现有的文档、网络研讨会，当然还有许多其他可用资源。

以下是启动最大尺寸的一维内核的合适方法吗？ GPU 上存在的所有 N 个线程都会执行这项工作吗？

这当然是可能的，所有启动的线程（假设它称为 N）都可以执行工作，并且它会启动一个最大 (1D) 大小的网格。但是你为什么要这样做呢？大多数 cuda 编程方法并不是以此为目标开始的。网格的大小应适合算法。如果 1D 网格大小似乎是一个限制器，您可以通过在内核中执行循环来解决每个线程处理多个数据元素的问题，或者启动 2D 网格以绕过 1D 网格限制。 cc3.x 设备的限制已经扩大。

与 cudaDeviceProp::maxThreadsDim[0] 相关的属性 cudaDeviceProp::maxThreadsPerBlock 是什么？它们有何不同？ cudaDeviceProp::maxThreadsPerBlock 可以代替上面的 cudaDeviceProp::maxThreadsDim[0] 吗？

首先是对多维块中线程总数的限制（即 threads_x*threads_y*threads_z）。第二个是对第一个维度 (x) 大小的限制。对于 1D 线程块，它们是可互换的，因为 y 和 z 维度为 1。对于多维块，存在多维限制以告知用户线程块，例如

maxThreadsDim[0]*maxThreadsDim[1]*maxThreadsDim[2]

是不合法的。

假设我想在块中的线程之间平均分配块的共享内存，并且我希望每个线程都可以使用最大数量的共享内存。那我应该最大化块的数量，并最小化每个块的线程数，对吧？

再次，我对方法论有点怀疑。但是，是的，每个线程可能的共享内存字节的理论最大值将通过线程数最少的线程块来实现。然而，允许一个线程块使用所有可用的共享内存可能会导致一次只有一个线程块可以驻留在一个 SM 上。这可能会对入住率产生负面影响，进而可能对性能产生负面影响。对于选择线程块大小以最大化性能，有许多有用的建议。我不能在这里总结它们。但是我们想选择线程块大小作为 warp 大小的倍数，我们通常希望每个线程块有多个 warp，在所有其他条件相同的情况下，我们希望启用最大占用率（这与可以驻留在一个线程块上的线程块数有关SM）。

只是为了确认（在查看有关 SO 的相关问题之后），在上面的线性（1D）网格/块方案中，全局唯一线程索引是 unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x。对吧？

是的，对于一维线程块和网格结构，这一行将给出一个全局唯一的线程ID：

unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;