DAY57:阅读Execution Configuration

B.21. Execution Configuration

Any call to a __global__ function must specify the execution configuration for that call. The execution configuration defines the dimension of the grid and blocks that will be used to execute the function on the device, as well as the associated stream (see CUDA C Runtime for a description of streams).

The execution configuration is specified by inserting an expression of the form <<< Dg, Db, Ns, S >>> between the function name and the parenthesized argument list, where:

• Dg is of type dim3 (see dim3) and specifies the dimension and size of the grid, such that Dg.x * Dg.y * Dg.z equals the number of blocks being launched;
• Db is of type dim3 (see dim3) and specifies the dimension and size of each block, such that Db.x * Db.y * Db.z equals the number of threads per block;
• Ns is of type size_t and specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory; this dynamically allocated memory is used by any of the variables declared as an external array as mentioned in __shared__; Ns is an optional argument which defaults to 0;
• S is of type cudaStream_t and specifies the associated stream; S is an optional argument which defaults to 0.

As an example, a function declared as

__global__ void Func(float* parameter);

must be called like this:

Func<<< Dg, Db, Ns >>>(parameter);

The arguments to the execution configuration are evaluated before the actual function arguments.

The function call will fail if Dg or Db are greater than the maximum sizes allowed for the device as specified in Compute Capabilities, or if Ns is greater than the maximum amount of shared memory available on the device, minus the amount of shared memory required for static allocation.

本文备注/经验分享：

这节讲的是, kernel的启动配置.

我们都知道, CUDA依靠海量线程并行, 来发挥它的高性能。之前曾经举过例子, 一块CPU, 可能几十个线程(逻辑核心), 双路的CPU系统, 往往不会超过100个核心.但一块普普通过的GTX1080, 就有20个SM, 而每个SM, 都能最多支持2048个线程. 这就4W多个线程能同时在执行了,远远不是CPU的规模能比的.所以在CPU上编程, 我们只需要进行: 函数名(参数);即可调用.而在GPU上, 我们却需要使用: kernel函数名<启动配置>(参数);才可以.这里的启动配置也即是本章节说的"执行配置".它有多种能配置的内容, 一共有4项, 其中的前两项则构成了海量线程的启动要求. 我们先来说说这个.如同你曾经总是能在CUDA中项目能看到的一样, 你会遭遇:

your_kernel<<<888, 666>>>(参数);

这里的888和666, 则是刚才说的启动配置的前两个项目: blocks数量, 和每个blocks里面的线程数量. 本例则是要求启动888个blocks, 每个blocks里面有666个线程.也就是一共会启动888 * 666, 大约50多W个线程.这里需要有几点注意的: (1)硬件本身(例如GTX1080)所能同时执行的线程数量有限, 例如最多4万个线程. 但是这里, 例如本例, 我们却可以直接要求50W个线程的启动规模. 这是因为, GPU会自动调度这些线程的. 总共50W个线程会分批发送到GPU的SM里执行的. 无需但心.

实际上, 计算能力相同的卡(例如GTX1080, 和300元的GTX1030, 计算能力都是6.1)，它们能同时通过启动配置要求的总线程规模是一样的.也就是无论实际的SM, SP有多少,它们的能启动的线程数的特性是一样的. 这个只取决于计算能力.但是GTX1080会执行的快很多. GT1030则会慢很多.这是因为对于同样的启动线程总数要求, 前者能同时执行的线程数量多, 后者能同时执行的数量少.前者可能分个几十个批次就全部自动上去执行完了.后者1030可能需要几百个批次.但是用户在写程序的时候不需要考虑这个,只要计算能力一样, <<<>>>里面的线程数量指定即可一样.

(2)实际应用中, 很多情况下, 是将问题切分给每个线程各自处理一部分的. 例如要绘制一块蛋糕, 每个线程可能只会绘制一个小部分, 假设某蛋糕项目, 需要100000个线程，但是你刚才看到, <<<>>>启动配置的前两项是分别指定了blocks数量和blocks里面的线程数量,此时往往会出现一种情况, 例如我使用1024个线程的blocks, 那么100000里面则需要有:100000 除以 1024, 得到97.6个blocks，请注意这里不能整除, 你无法直接要求0.6个blocks最后, 怎么办?此时应当直接启动98个blocks,然后在代码里面用if判断, 对于98个blocks * 1024线程/blocks = 100352, 大于100000个实际要求的额外352线程, 屏蔽掉它们.这样可以避免越界之类的问题.这是一个重要的需要注意的事项.大量kernel挂掉的问题和这个处理不当(边界)有关.

(3) 对于(2)中的blocks数量计算, 手册推荐的做法是: (a / b - 1) / b来计算, 其中a是总线程数量, b是每个blocks里面的线程数量.这个是对初学者的一个重要的迷惑地方.这个式子在干吗? 其实这个是一个常见的C整数运算的优化. 即快速向上取整. 它等价于:

count = a / b; if (a % b != 0) //最后一个blocks不能完整存在 { count++; }

这样就清晰很多(相比(a + b - 1) / b).我建议初学者先用这个式子, 而不要用手册上的. 避免困惑. (4)启动配置的前两项, 有时候不仅仅叫blocks数量和block里的线程数量配置,也叫"启动形状配置".这是因为往往很多问题, 需要使用2D或者3D形状的线程组合.例如block不仅仅是666个线程构成, 而是(16,16)这样的方形形状.同理grid不仅仅是888个blocks构成, 而是(1024, 1024)这样的方形的blocks.所以这就涉及到了线程索引和block索引的计算.建议用户自行看下blockIdx, threadIdx等等这些内置变量.这些内置变量论坛经常问到。 之前写过一个内置变量的章节, 有需要可以翻回去看看.这个是启动配置的前两个项目(启动形状)的解说.启动配置还有后两个项目, 一个是动态shared memory使用, 一个则是stream的使用.前者是一个非常重要的特性, 因为很多问题, 你无法直接在kernel里写成: __shared__ type var[N];的形式. 其中N是一个编译时刻的常量.而不能是变量.

如果N我无法提前确定好(在写代码的时候), 例如和具体问题的数据有关.那么则必须通过启动配置的第三个参数, 指定大小.之前论坛有一个PGI的客户, 因为错误的指定了这个大小为0, 导致kernel访存挂掉.所以需要注意.此外, 如果你的kernel部分使用的shared memory是固定大小的, 部分是不定的,则两者可以混合使用. 最终的实际大小是你这里的第三个参数指定的动态大小 + 你写死的静态大小.例如分别为10KB和8KB, 最终kernel执行的时候, 会使用18KB，这里需要注意.最后还需要注意的是, 所有计算能力, 除了7.0+的, 目前CUDA 9.2所支持的所有最大每个block能使用的shared memory大小(还记得shared memory是按照block进行资源分配, 而不是按照线程吗? 之前章节有说过), 均为48KB. 这个是个死的固定值.可以直接记住(已经从2010年到现在就没有变过了. 直到今天的7.0).而计算能力7.0如果想用更多的shared memory, 例如96KB, 则需要配合这里的启动参数配置的第三个动态参数才可以用到这么巨大. 总之手册上后面也又说过7.0使用巨大Shared memory的问题, 请看后面的该章节的如何启动配置即可.这里就不举例子了. 最后是Stream参数. 这个也很重要.但stream这个之前有过详细描述. 可以取得很多小kernel计算之间, 计算和传输之间, 同时进行.这也是一个非常重要的方面. 感兴趣的可以回头看看之前的章节单独说过这个. 此外, 需要补充的时,CUDA Driver API和OpenCL, 均没有这种方便易用的<<<>>>(也叫钻石形)启动配置.它们可以使用一个复杂的启动函数来启动kernel.<<<>>>这种语法, 大量的简化了用户启动kernel的负担(就像你启动一个普通的CPU上的: 函数(参数);一样简单, 只是多了4个值)，是CUDA Runtime API的特有福利. 因此你可以看到NV的软件环境有效的降低了你的开发成本,等效的增加了受益, 并延长了你的生命.所以N卡, 是从来不亏的选择.

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖