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

Question

这是一个关于如何确定CUDA网格,块和线程大小的问题.这是对此处发布的问题的另一个问题:

/sf/answers/395068691/

在此链接之后,talonmies的答案包含一个代码片段(见下文).我不理解评论"通常由调整和硬件约束选择的值".

我没有找到一个很好的解释或澄清,在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);

顺便说一句,我从上面的链接开始我的问题,因为它部分回答了我的第一个问题.如果这不是在Stack Overflow上提问的正确方法,请原谅或建议我.

Answer 1

答案有两个部分(我写了).一部分易于量化,另一部分更具经验性.

硬件限制:

这是容易量化的部分.当前CUDA编程指南的附录F列出了许多硬限制,这些限制限制了内核启动每个块的线程数.如果超过其中任何一个,您的内核将永远不会运行.它们大致可归纳为:

1. 每个块总共不能超过512/1024个线程(Compute Capability 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/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. 每个块的线程数应该是warp大小的四舍五入,在所有当前硬件上都是32.
2. GPU上的每个流式多处理器单元必须具有足够的活动warp,以充分隐藏架构的所有不同内存和指令流水线延迟,并实现最大吞吐量.这里的正统方法是尝试实现最佳硬件占用率(Roger Dahl的回答是指).

第二点是一个巨大的话题,我怀疑任何人都会尝试在一个StackOverflow答案中覆盖它.还有人写解决问题方面的定量分析博士论文(见本演示由加州大学伯克利分校的瓦西里·沃尔科夫和本文由亨利王从多伦多大学的真正的问题有多么复杂的例子).

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

Answer 2

上面的答案指出了块大小如何影响性能,并根据占用率最大化为其选择提出了一个通用的启发式方法.不希望提供选择块大小的标准,值得一提的是CUDA 6.5(现在在Release Candidate版本中)包括几个新的运行时函数,以帮助进行占用计算和启动配置,请参阅

CUDA Pro提示:占用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,因此可以cuOccupancyMaxPotentialBlockSize在驱动程序API代码中使用与上面示例中为运行时API显示的相同的方式.

Answer 3

通常选择块大小以最大化"占用".搜索CUDA占用情况以获取更多信息.特别是,请参阅CUDA占用计算器电子表格.