【问题标题】:How do I choose grid and block dimensions for CUDA kernels?如何为 CUDA 内核选择网格和块尺寸?
【发布时间】:2012-04-16 16:09:19
【问题描述】:

这是一个关于如何确定 CUDA 网格、块和线程大小的问题。这是here 发布的问题的附加问题。

在此链接之后,来自 talonmies 的答案包含一个代码 sn-p(见下文)。我不明白“通常由调整和硬件限制选择的值”的评论。

我在 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);

【问题讨论】:

    标签: performance optimization cuda gpu nvidia


    【解决方案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 的共享内存(计算 1.x/2.x-6.2/7.0)

    如果您保持在这些限制范围内,您可以成功编译的任何内核都将启动而不会出错。

    性能调优:

    这是经验部分。您在上述硬件限制内选择的每个块的线程数可以并且确实会影响在硬件上运行的代码的性能。每个代码的行为方式都会有所不同,量化它的唯一真正方法是仔细的基准测试和分析。但同样,非常粗略地总结:

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

    第二点是一个巨大的话题,我怀疑有人会尝试在一个 StackOverflow 答案中涵盖它。有人围绕问题的各个方面的定量分析撰写博士论文(参见加州大学伯克利分校的 Vasily Volkov 的 this presentation 和多伦多大学的 Henry Wong 的 this paper,以了解问题的复杂程度)。

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

    【讨论】:

    • "每个块的线程数必须是 warp 大小的整数倍" 这不是必须的,但如果不是必须的,你会浪费资源。我注意到 cudaGetLastError 在内核启动后返回的 cudaErrorInvalidValue 有太多块(看起来计算 2.0 无法处理 10 亿个块,计算 5.0 可以) - 所以这里也有限制。
    • 您的 Vasili Volkov 链接已失效。我假设你喜欢他 2010 年 9 月的文章:Better Performance at Lower Occupancy 文章(目前在 nvidia.com/content/gtc-2010/pdfs/2238_gtc2010.pdf 找到),这里有一个带有代码的 bitbucket:bitbucket.org/rvuduc/volkov-gtc10
    【解决方案2】:

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

    CUDA Pro Tip: Occupancy API Simplifies Launch 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");
    
    }
    

    编辑

    cudaOccupancyMaxPotentialBlockSizecuda_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 显示的方式相同。

    【讨论】:

    • 我有两个问题。首先,何时应该选择网格大小作为 minGridSize 而不是手动计算的 gridSize。其次,您提到“该函数提供的值可以用作手动优化启动参数的起点。”-您的意思是启动参数仍需要手动优化吗?
    • 有没有关于如何计算 2D/3D 块尺寸的指导?就我而言,我正在寻找 2D 块尺寸。是否只是计算 x 和 y 因子相乘得出原始块大小的情况?
    • @GrahamDawes this 可能感兴趣。
    【解决方案3】:

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

    【讨论】:

      猜你喜欢
      • 2011-08-14
      • 2018-05-29
      • 2015-09-17
      • 2013-08-25
      相关资源
      最近更新 更多