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

Question

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

问题 1： 以下是启动 最大 大小的一维内核的合适方法吗？ 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 并不是真正的教程服务。您应该利用现有的documentation、webinars，当然还有许多其他可用资源。

以下是启动最大尺寸的一维内核的合适方法吗？ 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 上。这可能对占用产生负面影响，而这可能对性能产生负面影响。有许多有用的建议来选择线程块大小，以最大限度地提高性能。我不能在这里总结它们。但是我们想选择线程块大小作为线程块大小的倍数，我们通常希望每个线程块有多个线程块，并且在所有其他条件相同的情况下，我们希望启用最大占用（这与可以驻留在一个线程块的数量有关） SM）。

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

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

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