★ FEATURED ARTICLE 搞CUDA最让人上火的事之一就是kernel写对了却在启动那一行卡住kernel???这三个问号到底填多少grid、block、thread这三层概念很多人看了教程知道大概但真到自己写的时候还是靠猜、靠试、靠复制别人的模板。我刚开始接触CUDA时也这样对着blockIdx.x、threadIdx.x发懵调一个越界错误调了一整晚最后发现是索引换算写错了。这篇东西就把grid、block、thread的设置这件事从头到尾捋一遍它们是什么、为什么是三层而不是一层、启动参数怎么写、索引怎么换算、block大小和grid大小到底怎么选、实测性能差距有多大。适合刚学CUDA想跑通第一个kernel的人也适合已经能跑但性能总是不对劲的开发者。1. 三层结构不是规矩多而是GPU硬件逼出来的很多CPU多线程编程的人第一次用CUDA都会觉得这三层结构很啰嗦直接开N个线程干活不行吗为什么非得分成grid和block还真不行这是由GPU的物理执行方式决定的。1.1 GPU到底怎么执行一堆线程GPU不是CPU那种几个强核并发跑几个线程的模式。一个GPU里有大量计算单元但执行的基本调度单位不是单个线程而是一组线程。这一组通常叫一个warp每个warp包含32个线程。硬件调度器会挑一个warp让这32个线程执行同一条指令只是各自处理不同的数据。这就是常说的SIMTsingle instruction, multiple threads风格。好处是控制逻辑非常省一个指令可以同时喂给32个线程单位时间能压榨出巨大的吞吐量坏处是你没法让单个线程有自己的程序计数器线程的灵活度远不如CPU线程。也就是说从硬件角度看GPU根本没有一个线程随便跑的概念它天生就是按批处理线程的。所以thread这个最小单元软件上存在硬件上却是凑够32个才一起上。这也解释了为什么下面反复强调block大小最好跟32扯上关系。1.2 block存在的意义协作和同步那为什么不能整个grid直接就是一大坨线程因为硬件不可能让成千上万个线程随便互相通信、同步。如果所有线程都在同一个集合里同步一次就要让所有线程停下来等你这代价太高了。于是GPU设计了一个中间层block。一个block内部的线程可以共享一块叫shared memory的存储空间也可以执行__syncthreads()做block内同步。但block和block之间是不保证执行顺序的也不能互相同步它们各自独立被调度到某个计算单元上。你可以把这个结构类比成一个公司grid是整个项目block是项目里的各个小组thread是小组里的成员。小组内部成员可以随时讨论、碰头shared memory、同步但小组之间互不干涉项目最终由所有小组拼起来。CPU里那种一个全局线程池随便通信的模型到了GPU这里被刻意拆掉了。1.3 这种分层带来的可扩展性一开始我也觉得这是限制后来才体会到这是GPU能通吃老卡、新卡的关键。因为block是独立调度的你的代码写完之后硬件会根据实际的计算单元数量把block均匀塞进去。同一个kernel跑在一台只有几个计算单元的卡和一台有几十个计算单元的卡上都能跑只是block被并行处理的程度不一样。也就是说grid到底有多大理论上不需要跟物理硬件一一对应你只需要保证整体工作量等于数据量就行。这个特性在后面设置grid大小时非常有用也是为什么可以放心用grid-stride loop那类写法的根本原因。2. 启动配置怎么写才不出错语法、上限和默认值摸清了硬件逻辑接下来就是最实际的启动配置代码怎么写、哪些数字绝对不能超。2.1 启动语法和dim3先看最基本的语法一个CUDA kernel的启动长这样kernelgridDim, blockDim, sharedMemBytes, stream(args...);gridDim一个dim3类型的变量表示grid里有多少个block可以是一维、二维、三维。blockDim表示每个block里有多少个thread同样可以是多维。sharedMemBytes动态共享内存字节数默认0没有特殊需求就留空。stream指定运行在哪个流上默认0表示默认流。dim3是CUDA提供的一个简单结构体有三个成员x、y、z。如果只关心一维可以直接写整数比如kernel10, 128(d_data);这表示grid有10个block每个block有128个线程总共1280个线程。但多维场景必须用dim3。比如要处理一张二维图像通常会把线程组织成二维的dim3 block(16, 16); dim3 grid((width 15) / 16, (height 15) / 16);。为什么二维场景要这么排一方面是索引算起来贴合数据行列结构另一方面是访存局部性和shared memory的协作模式更自然后面第三章会详细讲。2.2 必须背下来的硬件上限这块是踩坑重灾区很多为什么我的kernel启动就报错都是因为超了下面这些上限限制项典型上限说明每个block的最大线程数1024这是blockDim.x * blockDim.y * blockDim.z的乘积上限不是每个维度都1024blockDim各维度上限x、y、z 都是1024但三维乘积不能超过1024gridDim.x上限2^31-1约21亿新架构下x维度非常大gridDim.y/z上限65535很多老卡和文档里都写这个值warp大小32硬件执行批次不是软限制是硬规定一个block的最大线程数限制1024即便老文档里也写一致容易搞混的是二维block。有的新手以为blockDim(64, 64)是64×644096个线程直接报错。实际上一旦乘积超过1024启动时就可能失败或者行为未定义。所以如果你写dim3 block(32, 32)那每个block就是1024个线程这已经是上限了想要16×16这种常用配置那就是256个线程。还有gridDim.y和gridDim.z上限65535这在做二维或三维大规模问题时需要注意。如果你处理的图像高度超过65535个block那y方向就超了得考虑换一种分块策略。2.3 上限背后是寄存器、共享内存和warp这些上限不是随便拍脑袋拍出来的。一个block的线程数之所以限制在1024是因为GPU的计算单元SM内部资源是有限的寄存器文件就那么大shared memory就那么多活动线程数也有上限。block太大一个SM塞不下资源就被撑爆了。还有一个隐藏规则block里的线程数如果不是32的整数倍会浪费warp。比如block100硬件调度的时候会把它当成4个warp来处理100整除32是3余4最后一个warp里只有4个线程干活剩下28个位置空转。这不会报错但确实影响效率。所以常规建议都是block大小取32的整数倍64、128、256、512这类。3. 索引换算最容易被写错但又必须一次写对启动配置只是把有多少线程声明出来了线程之间怎么分工靠的是索引换算。这一节我觉得是整个CUDA入门里最容易出bug的地方值得单独拿出来反复说。3.1 一维换算一维场景最简单。假设你要给一个有n个元素的数组做加法启动配置是int threads 256; int blocks (n threads - 1) / threads; addKernelblocks, threads(d_a, d_b, d_out, n);kernel里每个线程拿到的全局索引是int idx blockIdx.x * blockDim.x threadIdx.x;为什么是这个公式想象block是一摞车blockIdx是第几辆车blockDim是每辆车有多少座位threadIdx是车里的第几个座位。全局座位号就是前面的车一共装了多少人blockIdx.x * blockDim.x加上我这辆车里的座位号threadIdx.x。注意blocks的计算(n threads - 1) / threads是向上取整。比如n1000, threads2561000/2563.9向上取整就是4个block总线程数1024比实际数据多24个。这24个多余的线程必须在kernel里挡掉不然就越界了。3.2 二维图像场景图像处理是CUDA最常见的应用场景之一而图像天然是二维的。如果非要用一维索引硬算也不是不行但代码很难看共享内存的tile方式也很难做。通常的做法是把block也排成二维__global__ void imgKernel(const float* input, float* output, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width y height) { int idx y * width x; output[idx] input[idx] * factor; } }启动时dim3 block(16, 16); dim3 grid((width block.x - 1) / block.x, (height block.y - 1) / block.y); imgKernelgrid, block(d_input, d_output, width, height);每个线程通过x blockIdx.x * blockDim.x threadIdx.x得到自己在图像里的列号y同理得到行号最后idx y * width x把二维坐标线性化成一维内存偏移。这里的顺序很关键图像内存通常按行存储所以是行号乘以宽度加列号写成x * height y就是错的。3.3 越界判断和网格跨度循环上面例子里的if (x width y height)就是越界判断。因为(width 15) / 16向上取整后grid的边缘block里必然有线程落在图像外面。除非你的数据精度是16的倍数否则这个判断必须有。还有一类更通用的写法叫grid-stride loop特别适合数据量远大于线程总数的场景__global__ void addKernel(const float* a, const float* b, float* out, int n) { int stride gridDim.x * blockDim.x; for (int idx blockIdx.x * blockDim.x threadIdx.x; idx n; idx stride) { out[idx] a[idx] b[idx]; } }启动时grid可以设得比较固定比如等于某个合理值然后每个线程用步长循环处理多个元素。这个模式的好处是线程总数和数据量解耦了启动配置和算法逻辑更干净同时循环里连续访问idx的元素访存模式依然是合并访问性能不差。3.4 常见的索引错误我见过、也踩过几类典型错误一是把blockIdx.x和threadIdx.x搞反算出来的idx直接乱了。二是忘记边界判断数据量不是block的整数倍时越界读写程序可能不报错但算出错结果或者偶尔崩溃极难排查。三是误以为blockIdx.x全局唯一、threadIdx.x全局唯一其实threadIdx在每个block内部都是从头数的必须组合起来才是全局位置。四是二维场景把宽高搞混导致访问模式变成列优先性能急剧下降。排查这类问题的笨办法就是在kernel里先不加业务逻辑直接把idx写到输出数组里然后在CPU端验证一遍确认每个元素的索引映射正确了再往下加数据运算。这一步很值得做因为索引错了后面的所有优化都是在错误地基上盖楼。4. block和grid怎么定从能跑到跑得快配置能跑起来之后下一步就是性能。block大小和grid大小怎么选网上说法很多但核心逻辑其实可以拆成三步走。4.1 第一步确定block size的经验值初次调优我建议直接从256开始。理由很朴素256是32的8倍warp利用率完整256个线程对应活跃线程数不会太小能较好占满SM对大多数kernel256不会把寄存器或shared memory压到极限留出了调整余地。有些人喜欢用128有些人用512。实测下来在大多数简单计算型kernel里128、256、512的差距并不大真正拉开差距的是资源限制和访存模式。所以别在第一步上过度纠结先定一个256的默认值跑通再用下面的方法验证。也不要盲目上1024。1024是上限但寄存器、shared memory占用率高时一个1024线程的block可能导致SM只能驻留一个block反而带来调度和同步上的劣势。block大小不是单纯越大越好要看资源账。4.2 第二步用占用率检查资源瓶颈占用率occupancy指的是一个SM上实际活动的warp数除以理论最大warp数。这个指标大致反映了计算单元被填满的程度。你可以用性能分析工具看也可以自己大致估算。举一个简单例子。某GPU每个SM最多2048个线程寄存器文件65536个。假设你的kernel每个线程用32个寄存器block256一个block需要256×328192个寄存器每个SM最多同时容纳65536/81928个block共2048线程达到上限。再假设每个线程用64个寄存器一个block需要256×6416384个寄存器只能放65536/163844个block共1024线程占用率腰斩。这种时候把block降到128反而能放8个block共1024线程。虽然总占用率没变但block数量多了调度灵活度更高有时候性能还更好。所以block size的选择要结合每个kernel的资源开销来看不是一个固定值走天下。shared memory也是同样的逻辑。一个block能用8KB shared memorySM总共96KB如果block256每个block用8KB那最多驻留几个block是有限的。遇到这种kernel要么减小block要么减少shared memory使用量。4.3 第三步决定grid大小和负载均衡确定block之后grid大小通常就是总数据量 / block大小向上取整。这个没太多花活但要留意负载均衡。如果数据量很小比如只有1000个元素block256那就只要4个block。如果GPU有几十个SM4个block根本填不满大量计算单元空转。这种情况下可以换一个思路要么数据量太小不值得上GPU要么用grid-stride loop把grid大小设成SM数乘以每个SM能驻留的block数让每个线程处理少量元素反而能把整个GPU铺满。反过来如果数据量巨大而block设得太小比如block32grid可能有几百万个block虽然也能跑但每次block调度都有开销性能也会受影响。常规原则是在资源允许的前提下block尽量选大一些128及以上让grid数量可控。4.4 针对不同访存模式的微调上面是通用流程实际项目里还要看kernel的访存特征。如果是纯计算密集、每个线程独立处理一个元素的简单kernel256或512通常都不错重点看寄存器别超额。如果是访存密集型kernel比如矩阵运算、图像卷积block的形状常常取(16,16)或(32,8)这种矩形。因为这类算法往往要配合shared memory做tile缓存矩形block能更好地贴合数据块。比如做卷积每条线程要读周围邻域的像素把一个16×16线程块映射到16×16像素块shared memory里再加载一圈halo数据这样访存的重复读会少很多。如果是规约reduction这类需要block内协作的kernelblock取256、512比较常见。block太小则分块很多最后结果合并开销大block太大则同步和shared memory开销上升。而且这类kernel还有一个老生常谈的点最后一次block内规约时线程数按warp粒度折半甚至直接用shuffle指令这已经是kernel内部写法的事了但跟block大小也有关联。5. 实测数据与调优记录一个数组求和kernel的配置对比讲完理论上点实测数据给大家一个体感。我最近用一个很简单的数组求和kernel做了一组对比测试不同block大小下的耗时差异比想象中明显。5.1 测试环境与方法测试环境是某块中端独显数据量是1000万浮点数kernel把两个数组逐元素相加。为了排除启动开销干扰我先跑了一次warmup然后用CUDA事件cudaEventRecord计算实际kernel执行时间每组跑10次取中位数。kernel本身非常简单没有shared memory每条线程读两个数写一个数重点就是看不同启动配置对执行效率的影响。5.2 实测结果block大小grid大小耗时微秒备注32312500约 45网格过大调度开销明显64156250约 32略好仍偏高12878125约 24进入正常区间25639063约 21本次最优之一51219532约 22和256相当10249766约 24资源调度变粗略降从这个结果能看出block从32提到256性能从45微秒降到21微秒差距超过一倍。而256到1024之间差距很小甚至1024略回落。这就是前面说的不是越大越好的实证。不同GPU、不同驱动、不同数据类型肯定会有差异但趋势是一致的block过小时的调度开销和block过大时的资源僵硬都不如中间区间。5.3 调优过程里踩过的坑第一坑向上取整写错。(n threads - 1) / threads在整数运算下是正确的可有人图省事直接写n / threads数据量恰好整除时没事不整除时最后一个block覆盖不到尾部数据结果数组末尾永远是0。这种错非常隐蔽因为程序不崩溃数值也合理只是不对。第二坑以为1024一定最快。上面数据已经说了反而可能略慢而且一旦kernel的寄存器或shared memory占用上来1024会让占用率非常难看。正确的顺序是先资源估算再实测。第三坑改配置后没清错误。启动后的错误是异步的最好在kernel启动后立刻调用cudaGetLastError()检查一下或者在调试阶段用工具跑一遍。不然有时启动参数超限了程序表面正常结果全错排查起来怀疑人生。第四坑性能对比时没保持变量单一。有时候我一下改block一下改grid一下又改了kernel内循环结果性能变了不知道是谁的功劳。正确的做法是每次只改一个变量其他全部固定记录数据再比较。5.4 怎么快速定位启动配置导致的问题如果真的怀疑是启动配置引起的错误可以按下面的流程排查。第一件事用工具做内存和越界检查。老的工具叫cuda-memcheck新的工具叫compute-sanitizer都能帮你抓越界访问、未初始化内存这类问题。报出来的错误栈一般会告诉你哪个kernel、哪个索引出了问题。第二件事看启动配置是否超限。最简单的检查方式就是cudaGetLastError()它会立刻把上一次启动的错误码返回给你。如果是invalid configuration多半是block线程数乘积超了1024或者grid维度超了上限。第三件事分析占用率和资源瓶颈。用GPU厂商的profiling工具打开kernel分析看每个SM的理论占用率、实际占用率、寄存器溢出、shared memory bank冲突。这些信息比盲目试block size有效得多。通常走到第三步问题基本已经浮出水面了。要么是block过大导致占用率上不去要么是shared memory或寄存器超限导致编译器不得不做local memory溢出性能暴跌。这时候再回头调里的数字就是有依据的调而不是瞎调。6. 一点个人经验收尾写了这么多最后分享几条我觉得对新手最有用的经验。第一先把索引换算是基本功。启动配置再花哨索引错了就是零分。我习惯任何kernel都是先写一个只回写索引的调试版CPU端验证一次再开始写真正的业务逻辑。第二默认配置先从block256开始。它能覆盖绝大多数第一版kernel跑通之后再用分析工具看瓶颈决定调不调。没必要第一版就追求理论最优配置优化是迭代出来的不是拍脑袋出来的。第三每次只改一个变量。记录下block128、256、512、1024各组的耗时心里有个数据表你调优的方向感会清楚很多。调过一次之后你会发现自己对该往哪个方向调的判断力会明显提升。
阅读完成 · 觉得有帮助？
