如何为 CUDA 内核选择网格和块尺寸？

这是一个关于如何确定 CUDA 网格、块和线程大小的问题。这是一个额外的问题，一张张贴 给你。

在这个链接之后，talonmies 的答案包含一个代码片段(见下文)。我不理解“通常由调优和硬件约束选择的值”这样的评论。

我还没有找到一个很好的解释或澄清，解释这在 CUDA 文档。总之，我的问题是如何确定给定以下代码的最佳 blocksize(线程数) :

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);
133446 次浏览
小开

块大小通常是选择最大化的“占用”。搜索 CUDA 占用更多信息。特别是，请参阅 CUDA 占用率计算器电子表格。

小开
最佳答案

这个答案有两部分(我写的)。一部分是容易量化的，另一部分是更实证的。

硬件限制:

这是容易量化的部分。当前 CUDA 编程指南的附录 F 列出了一些硬性限制，这些限制限制了内核启动时每个块可以有多少线程。如果超过其中任何一个，您的内核将永远不会运行。它们大致可概括如下:

  1. 每个块的线程总数不能超过512/1024个(分别为 计算能力1.x 或2.x 及更高版本)
  2. 每个块的最大尺寸限制为 [512,512,64]/[1024,1024,64](计算1.x/2.x 或更高版本)
  3. 每个块不能使用超过8k/16k/32k/64k/32k/64k/32k/64k/32k/32k/32k/64k 寄存器总数 (计算1.0,1.1/1.2,1.3.2. x-/3.0/3.2/3.5-5.2/5.3/6-6.1/6.2/7.0)
  4. 每个块不能消耗超过16kb/48kb/96kb 的共享内存(Compute 1. x/2.x-6.2/7.0)

如果您保持在这些限制范围内，那么您可以成功编译的任何内核都将正常启动。

性能调优

这是经验部分。在上面列出的硬件约束中选择的每个块的线程数可以并且确实会影响在硬件上运行的代码的性能。每个代码的行为方式将是不同的，量化它的唯一真正方法是通过仔细的基准测试和概要分析。但是，再次，非常粗略地总结:

  1. 每个块的线程数应该是经纱大小的整数倍，在所有当前硬件上是32。
  2. 图形处理器上的每个流式多处理器单元必须有足够的活动翘曲，以充分隐藏架构中所有不同的内存和指令管线化延迟，并实现最大吞吐量。这里的正统方法是尝试实现最佳硬件占用率(Roger Dahl 的回答指的是)。

第二点是一个巨大的话题，我怀疑是否有人会尝试用一个 StackOverflow 答案来涵盖它。有人写博士论文围绕着问题的定量分析(见加州大学伯克利分校的瓦西里 · 沃尔科夫的 这个展示和多伦多大学的亨利 · 王的 这张纸，他们举例说明这个问题到底有多复杂)。

在入门级别，您应该主要意识到您选择的块大小(在上述约束定义的合法块大小范围内)可以并且确实影响代码的运行速度，但这取决于您拥有的硬件和正在运行的代码。通过基准测试，您可能会发现大多数非平凡的代码在每个块范围的128-512个线程中都有一个“最佳位置”，但是这需要您进行一些分析来找到这个位置。好消息是，因为您的工作范围是经纱大小的数倍，所以搜索空间非常有限，对于给定的代码片段，最佳配置相对容易找到。

小开

上面的答案指出了块大小如何影响性能，并提出了一个基于占用率最大化的选择的通用启发式算法。不想提供 的标准来选择块大小，值得一提的是 CUDA 6.5(现在的候选发布版本)包含了几个新的运行时函数来帮助计算占用率和启动配置，参见

CUDA 专业提示: 占用 API 简化启动配置

其中一个有用的函数是 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开始，人们需要根据 API 建议的1D 块大小计算自己的2D/3D 块尺寸。

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


提交答案