如何为CUDA内核select网格和块维度？

这是一个关于如何确定CUDA网格，块和线程大小的问题。 这是在这里发布的另一个问题：

https://stackoverflow.com/a/5643838/1292251

在这个链接之后，来自talonmies的答案包含一个代码片段（见下文）。 我不明白“调整和硬件限制通常select的价值”的评论。

我还没有find一个很好的解释或澄清，在CUDA文档中解释这一点。 总之，我的问题是如何确定给出以下代码的最佳块大小 （=线程数）：

const int n = 128 * 1024; int blocksize = 512; // value usually chosen by tuning and hardware constraints int nblocks = n / nthreads; // value determine by block size and total work madd<<<nblocks,blocksize>>>mAdd(A,B,C,n); 

顺便说一句，我开始我的问题与上面的链接，因为它部分回答我的第一个问题。 如果这不是一个正确的方式提出堆栈溢出的问题，请原谅或劝告我。

Answer 1

这个答案有两个部分（我写的）。 一部分容易量化，另一部分更具经验性。

硬件限制：

这是容易量化的部分。 当前CUDA编程指南的附录F列出了一些限制内核启动每块可以有多less个线程的硬限制。 如果你超过了这些，你的内核永远不会运行。 他们大致可以概括为：

1. 每个块的总数不能超过512/1024（分别为Compute Capability 1.x或2.x-3.x）
2. 每个块的最大尺寸限制为[512,512,64] / [1024,1024,64]（Compute 1.x / 2.x）
3. 每块总共不能消耗超过8k / 16k / 32k的寄存器（Compute 1.0,1.1 / 1.2,1.3 / 2.x）
4. 每块不能占用16kb / 48kb以上的共享内存（Compute 1.x / 2.x）

如果你保持在这个限制之内，你可以成功编译的任何内核都将无误地启动。

性能调整：

这是实证部分。 在上面列出的硬件约束条件下，您select的每个块的线程数量可能会影响在硬件上运行的代码的性能。 每个代码的行为将是不同的，唯一真正的量化方法是通过仔细的基准testing和分析。 但是，粗略地总结一下：

1. 每个块的线程数应该是warp大小的整数倍，在当前所有硬件上都是32。
2. GPU上的每个stream式多处理器单元必须具有足够的活动warp来充分隐藏架构的所有不同内存和指令stream水线延迟，并实现最大吞吐量。 这里的正统方法是尝试实现最佳的硬件占用（ Roger Dahl的答案是指的是）。

第二点是一个巨大的话题，我怀疑有人会尝试和覆盖在一个单一的StackOverflow答案。 有些人围绕对问题的定量分析写博士论文（参见加州大学伯克利分校的瓦西里·沃尔科夫（Vasily Volkov）和多伦多大学的亨利·王（Henry Wong）的这篇论文 ，来说明这个问题究竟有多复杂）。

在入门级，你应该主要知道你select的块大小（在由上面的约束定义的合法块大小的范围内）可以并且对代码的运行速度有影响，但是它取决于硬件你有和你正在运行的代码。 通过基准testing，您可能会发现大多数不重要的代码在每个块范围128-512个线程中都有一个“最佳位置”，但是您需要进行一些分析才能find它的位置。 好消息是，因为您正在使用多倍的warp尺寸，所以search空间是非常有限的，对于给定的一段代码来说，最好的configuration相对容易find。

Answer 2

上面的答案指出了区块大小如何影响性能，并提出了一个基于占用率最大化的常见启发式select。 不需要提供select块大小的标准，值得一提的是，CUDA 6.5（现在处于Release Candidate版本）包含几个新的运行时function，以帮助进行占用计算和启动configuration，请参阅

CUDA Pro提示：占用API简化了启动configuration

其中一个有用的函数是cudaOccupancyMaxPotentialBlockSize ，它启发式地计算可达到最大占用率的块大小。 然后可以使用该函数提供的值作为手动优化启动参数的起点。 下面是一个小例子。

 #include <stdio.h> /************************/ /* TEST KERNEL FUNCTION */ /************************/ __global__ void MyKernel(int *a, int *b, int *c, int N) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if (idx < N) { c[idx] = a[idx] + b[idx]; } } /********/ /* MAIN */ /********/ void main() { const int N = 1000000; int blockSize; // The launch configurator returned block size int minGridSize; // The minimum grid size needed to achieve the maximum occupancy for a full device launch int gridSize; // The actual grid size needed, based on input size int* h_vec1 = (int*) malloc(N*sizeof(int)); int* h_vec2 = (int*) malloc(N*sizeof(int)); int* h_vec3 = (int*) malloc(N*sizeof(int)); int* h_vec4 = (int*) malloc(N*sizeof(int)); int* d_vec1; cudaMalloc((void**)&d_vec1, N*sizeof(int)); int* d_vec2; cudaMalloc((void**)&d_vec2, N*sizeof(int)); int* d_vec3; cudaMalloc((void**)&d_vec3, N*sizeof(int)); for (int i=0; i<N; i++) { h_vec1[i] = 10; h_vec2[i] = 20; h_vec4[i] = h_vec1[i] + h_vec2[i]; } cudaMemcpy(d_vec1, h_vec1, N*sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(d_vec2, h_vec2, N*sizeof(int), cudaMemcpyHostToDevice); float time; cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, 0); cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, MyKernel, 0, N); // Round up according to array size gridSize = (N + blockSize - 1) / blockSize; cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&time, start, stop); printf("Occupancy calculator elapsed time: %3.3f ms \n", time); cudaEventRecord(start, 0); MyKernel<<<gridSize, blockSize>>>(d_vec1, d_vec2, d_vec3, N); cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&time, start, stop); printf("Kernel elapsed time: %3.3f ms \n", time); printf("Blocksize %i\n", blockSize); cudaMemcpy(h_vec3, d_vec3, N*sizeof(int), cudaMemcpyDeviceToHost); for (int i=0; i<N; i++) { if (h_vec3[i] != h_vec4[i]) { printf("Error at i = %i! Host = %i; Device = %i\n", i, h_vec4[i], h_vec3[i]); return; }; } printf("Test passed\n"); } 

编辑

cudaOccupancyMaxPotentialBlockSize在cuda_runtime.h文件中定义，其定义如下：

 template<class T> __inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSize( int *minGridSize, int *blockSize, T func, size_t dynamicSMemSize = 0, int blockSizeLimit = 0) { return cudaOccupancyMaxPotentialBlockSizeVariableSMem(minGridSize, blockSize, func, __cudaOccupancyB2DHelper(dynamicSMemSize), blockSizeLimit); } 

参数的含义如下

 minGridSize = Suggested min grid size to achieve a full machine launch. blockSize = Suggested block size to achieve maximum occupancy. func = Kernel function. dynamicSMemSize = Size of dynamically allocated shared memory. Of course, it is known at runtime before any kernel launch. The size of the statically allocated shared memory is not needed as it is inferred by the properties of func. blockSizeLimit = Maximum size for each block. In the case of 1D kernels, it can coincide with the number of input elements. 

请注意，从CUDA 6.5开始，需要根据APIbuild议的一维块大小计算自己的2D / 3D块大小。

另请注意，CUDA驱动程序API包含用于占用率计算的function上等效的API，因此可以在驱动程序API代码中使用cuOccupancyMaxPotentialBlockSize ，与上例中的运行时API所示的方式相同。

Answer 3

通常select块大小来最大化“占用”。 searchCUDA入住以了解更多信息。 具体来说，请参阅CUDA占用率计算器电子表格。