如何选择CUDA核心的网格和块维度?

139
这是一个关于如何确定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);
3个回答

176

这个问题有两个部分(由我编写)。其中一个部分很容易量化,另一个则更具经验性。

硬件约束:

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

  1. 每个块总共不能超过512/1024个线程(计算能力1.x或2.x及更高版本)。
  2. 每个块的最大尺寸限制为[512,512,64]/[1024,1024,64](Compute 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的共享内存(计算1.x/2.x-6.2/7.0)。

如果您在这些限制内,任何一个您可以成功编译的内核都会成功启动,没有错误。

性能调优:

这是经验性的部分。您在硬件约束范围内选择的每个块的线程数量可以并且确实会影响代码在硬件上运行时的性能。每个代码的行为都不同,唯一真正量化它的方法是通过仔细的基准测试和分析。但是,大致概括:

  1. 每个块的线程数应该是warp大小的整数倍,当前硬件上warp大小为32。
  • GPU上每个流多处理器单元必须有足够的活跃线程束来充分隐藏架构中不同内存和指令流水线延迟,并达到最大吞吐量。正统的方法是尝试实现最佳硬件占用率(即 Roger Dahl's answer所指的内容)。
  • 第二点是一个非常庞大的课题,我怀疑没有人会试图在单个StackOverflow答案中全面覆盖它。有人写出了关于该问题方面的量化分析的博士论文(例如,参见来自加州大学伯克利分校的Vasily Volkov的这个演示文稿和来自多伦多大学的Henry Wong的这篇论文)。

    在入门级别上,您应该意识到您选择的块大小(在上述约束条件定义的合法块大小范围内)可能会对代码运行速度产生影响,但这取决于您拥有的硬件和正在运行的代码。通过基准测试,您可能会发现大多数复杂代码在每个块128-512个线程之间有一个“甜点”,但需要您进行一些分析来确定甜点的位置。好消息是,因为您正在使用线程束的倍数,搜索空间非常有限,对于给定代码的最佳配置相对容易找到。


    2
    每个块的线程数必须是warp大小的整数倍,否则会浪费资源。我注意到,在使用太多块进行内核启动后,cudaGetLastError会返回cudaErrorInvalidValue(看起来compute 2.0无法处理10亿个块,而compute 5.0可以)--因此这里也有限制。 - masterxilo
    4
    您提供的 Vasili Volkov 链接已失效。我猜您是想链接到他在2010年9月发布的“更低占用率下提高性能”的文章(现在可以在http://www.nvidia.com/content/gtc-2010/pdfs/2238_gtc2010.pdf找到),这里有一个带有代码的 bitbucket 链接:https://bitbucket.org/rvuduc/volkov-gtc10。 - ofer.sheffer

    47
    以上回答指出了块大小如何影响性能,并建议基于占用率最大化的常见启发式方法来选择块大小。并不想提供选择块大小的标准,值得一提的是,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代码中使用cuOccupancyMaxPotentialBlockSize,方法与上面示例中运行时API相同。


    3
    我有两个问题。首先,何时应该选择minGridSize作为手动计算的gridSize。其次,您提到:“该函数提供的值可以用作手动优化启动参数的起点。”- 您是否意味着仍需要手动优化启动参数? - nurabha
    有没有关于如何计算2D/3D块尺寸的指导?在我的情况下,我正在寻找2D块尺寸。这只是计算x和y因子,当它们相乘时给出原始块大小的情况吗? - Graham Dawes
    1
    @GrahamDawes 这个链接可能会对你有所帮助:https://dev59.com/JpDea4cB1Zd3GeqPZkp3#33247118 - Robert Crovella

    9

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


    网页内容由stack overflow 提供, 点击上面的
    可以查看英文原文,
    原文链接