在CUDA里写kernel，最难往往不是把算法写对，而是把启动参数写对。grid、block、thread这三个词，背概念的时候好像都懂，一旦自己动手配kernel<<<grid, block>>>，就会发现各种问题：block到底设多大？grid怎么算？二维三维怎么排？为什么换个block大小性能差那么多？

这篇文章我把这些年调参数的经验掰开揉碎讲清楚，从底层的线程组织逻辑、硬件限制，到一维和二维任务的配置方法，再到常见启动报错和性能排查思路，都会覆盖到。刚入门CUDA的同学可以直接照着改；写过一段时间kernel但总觉得性能不对的同学，也能从里面找到不少之前没注意过的细节。

1. 先搞懂CUDA的线程层级：为什么是grid、block、thread三层

1.1 从硬件视角看三层结构

我在刚学CUDA的时候，最大的困惑就是：既然要开一堆线程来处理数据，为什么不直接给一个“线程总数”，比如开100万个线程，非要搞出grid、block两层？直接kernel<<<1000000>>>不就行了？

等接触了实际硬件之后才明白，这个设计不是多此一举，而是GPU架构决定的。

一颗GPU上有很多个SM（Streaming Multiprocessor，流式多处理器），每个SM里面又有一组CUDA core。你在host端启动kernel时，发给GPU的是一个由block组成的grid；GPU的后端调度器会把一个个block分配到不同的SM上去执行。一个block一旦被指派给某个SM，只会待在那个SM上，不会中途跑到别的SM。

这带来的第一个好处是硬件异构性下的代码可移植性。不同GPU的SM数量差很多，低端卡可能只有几组SM，高端卡有几十上百组SM。如果你显式地指定“100万个线程”，那不同显卡怎么分配？但如果你的代码写的是“这个算式的每个输出点由一个线程负责，每个block放256个线程，grid放3907个block”，那不管GPU有几个SM，调度器都能按比例把block摊到所有SM上。写一次，在不同配置的卡上都能跑。

第二个好处是协作边界清晰。一个block内的线程可以通过共享内存（shared memory）交换数据，可以通过__syncthreads()做同步，这种协作效率极高，相当于同一小组内互相配合。但不同block之间没有这种直接通信手段，想跨block协作就得靠全局内存、原子操作或者干脆重新启一个kernel。这种设计有点像把任务拆成一个个可以独立推进的工作包，而GPU的任务调度器只需要在block粒度上做分配，不需要管线程和线程之间的复杂依赖。

第三个好处是天然适配多维数据。grid、block本身都是三维的dim3类型，可以很方便地把像素坐标、体素坐标映射成线程坐标。处理一张1024x768的图像，直接用block(16, 16)，再算一下grid尺寸就行，比手动把二维坐标压成一维再逆推回去直观得多。

1.2 warp：真正的最小调度单位

很多初学者以为GPU调度是以thread为单位，这是个常见误解。真正的调度单位是warp——一组32个线程。

SM执行指令时，不会一个线程一个线程地派发，而是一次拉一个warp的32个线程，让它们在同一时刻执行同一条指令。这是SIMT（Single Instruction, Multiple Thread）模型，形式上有点像CPU的SIMD，但区别在于warp内每个线程可以有自己的数据、自己的寄存器，只是“指挥”是共享的。

理解这一点，对设置block非常关键。blockDim.x最好取32的倍数，就是为了保证每个warp都被填满。假如你设block为48个线程，硬件会把它切成两个warp：第一个warp 32个线程满编，第二个warp只有16个线程。当第二个warp被调度时，有一半的lane是空闲的，白白浪费指令周期的执行能力。

这个细节不直接影响正确性，但会直接反映在性能上。所以GPU上跑kernel，最常见的block配置基本都是128、256、512这类32的倍数。这不是什么玄学规则，单纯是warp分组决定的。

另外一个和warp相关的现象叫分支分歧（divergence）。如果一个warp内的一部分线程走if分支，另一部分线程走else分支，硬件没法让32个线程同时执行两条路径，只能先执行if路径（另一半线程被屏蔽），再执行else路径（另一半线程被屏蔽），最终耗时相当于两条路径相加。这是SIMT模型固有的代价，在设置block和规划数据映射时最好尽量避开跨warp的分支。

2. block内部的组织方式与线程索引计算

2.1 内置变量到底怎么用

CUDA提供了一组内置变量，用来在kernel里定位当前线程的身份：

• threadIdx：dim3类型，表示当前线程在它所在block内的坐标。
• blockIdx：dim3类型，表示当前block在整个grid内的坐标。
• blockDim：dim3类型，表示block的尺寸，也就是每个维度上放了多少个线程。
• gridDim：dim3类型，表示grid的尺寸，也就是每个维度上放了多少个block。

这组变量在kernel里可以直接读，不需要自己声明，也不需要传入。

一维任务里，把线程映射到数组下标是最常见的操作：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式可以理解为：当前线程在整个grid里的“绝对位置” = 前面block里所有线程 + 自己block内的线程位置。你可以手算验证一下：假设blockDim.x = 256，blockIdx.x = 3，threadIdx.x = 42，那么tid = 3 * 256 + 42 = 810。

二维任务里，公式会变成两个方向各算一次：

int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; int idx = row * width + col; // 二维坐标压回一维下标

这里最容易搞混的是row和col的对应。row由y方向决定，col由x方向决定，因为内存里一行是连续摆放的，所以col对应连续内存方向。如果把x和y记反了，程序也能跑，但内存访问会变成跳着访问，性能损失非常明显。

2.2 一维任务为什么默认x方向

处理一维数组时，90%的人都只用到x方向，即gridDim.x和blockDim.x。这是因为x方向有硬件层面的特殊地位：

• 线程块内的x维度是内存访问最连续的方向，合并访问（coalesced access）最容易实现。
• grid在x方向的上限最大，x方向最多可以放2^31 - 1个block，远超y和z方向的65535上限。

所以即使你的业务可以拆成二维，如果数据和计算本质上是一维的，就老老实实用x方向。

有人会问，那数据超过x方向上限怎么办？答案是加grid-stride loop，让一个线程循环处理多个元素，而不是拼命拉大grid。比如n是1000万：

__global__ void vectorAddStride(const float* A, const float* B, float* C, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; int stride = gridDim.x * blockDim.x; for (int i = idx; i < n; i += stride) { C[i] = A[i] + B[i]; } }

这里stride等于整个grid能承载的总线程数。你定义一个差不多的grid大小，比如几万个线程，然后循环步进，直到覆盖完所有数据。

grid-stride loop还有一个额外好处：grid大小可以和设备SM数量挂钩，比如取的block总数是SM数量的几倍，让负载更均衡，而不用每次根据数据大小临时算grid。这个模式在处理超大数组时几乎是标准写法。

2.3 二维和三维block与数据排布

图像处理是典型的二维数据场景。假设一张图宽W、高H，想让每个线程处理一个像素，就可以用二维block：

dim3 block(16, 16); dim3 grid((W + 15) / 16, (H + 15) / 16);

每个block内是16x16的线程矩阵，同一block覆盖原图一个16x16的像素块。相邻线程访问相邻像素，在x方向上是连续合并的，所以这个方案既直观又在访存上占优。

三维block的写法也一样，只是多了一个z方向。但三维block里有个很多人踩过的坑：blockDim.z的上限只有64，不是1024。如果你把z方向设成128或者256，kernel启动直接报错。三维体数据处理通常建议把z方向切成小块，或者干脆把数据按平面的方式组织成多个二维block。

无论用几维block，边界处理都是必须做的。比如图像宽度W不是16的倍数，那grid的边缘block里就有部分线程越界。kernel里必须有类似if (row < H && col < W)的判断，否则这些线程会访问非法内存，轻则算错数据，重则崩溃。

3. 配置参数怎么定：硬件限制、占用率与性能权衡

3.1 硬件限制速查表

拿到一张卡，首先要搞清楚它的限制。我整理了一张常见的上限表，适用计算能力3.0及以上的大多数GPU：

重点强调：

第一个是blockDim.z上限64，这是一个很容易被遗忘的限制。体数据处理的时候，有人把block设成(16, 16, 8)，正好是2048线程，还超过1024上限，直接非法。三维block设计时一定要先查这个表。

第二个是每个SM最大block数。即使每个block线程数很小，SM能容纳的block数也是有限的。比如某个架构每个SM最多32个block，如果你的block只放64个线程，那最多驻留32 x 64 = 2048线程。但如果block是32线程，32个block只能放1024线程，SM的2048线程能力根本用不满。所以block太小反而会拖累占用率。

3.2 配置选择的核心原则

配置顺序永远是：先定block，再算grid。grid大小通常由数据规模和block大小做除法得到，不存在什么“最佳grid”的说法——它是一个依赖变量。

block大小的选择主要看这几个因素：

第一，必须是32的倍数。前面已经解释过warp对齐的原因。除非你有非常特殊的理由，否则不要打破这条。

第二，推荐区间在128到512之间。block太小，比如32或64，会让每个SM承载的block数过多，可能触到SM的最大block限制，而且调度开销占比相对变大。block太大，比如直接上1024，也不是不行，但每个SM只能驻留两个block。如果一个block在等__syncthreads()，另一个block也在等，SM就空了。小块block能在同一个SM上放更多独立工作包，如果一个block因为访存或同步暂停，调度器可以立刻切到另一个block的warp继续执行，这是隐藏延迟的关键机制。

第三，根据任务目标合理匹配维度。一维任务用x方向，二维任务用二维block，不要强行把一维任务映射到二维block上。虽然索引公式能算，但你牺牲了代码可读性，还可能引入不必要的非连续访问。

我来举个直观的例子。数据有100万个元素，选block=256，那grid就是ceil(1000000 / 256) = 3907。尾数不是整除，最后一个block里有896个线程在空转或等待边界判断。如果选block=128，grid=7813，空转的线程规模会更细。这种差异对性能影响不大，关键还是让总线程数覆盖任务，并且加好边界判断。

3.3 占用率计算与“100%”的迷信

占用率（occupancy）指的是SM上活跃warp数和SM能容纳的最大warp数的比值。很多初学者一听到“占用率越高越好”，就疯狂调小block，这个方向其实是错的。

举个例子。某个现代架构每个SM最多2048个线程，也就是64个warp（每个warp 32线程）。如果你用block=256，那无论block内有没有别的资源限制，SM最多能驻留8个block（2048/256=8），占用率是100%。block=128，驻留16个block，也是100%。block=512，驻留4个block，还是100%。只有block小到一定程度，比如block=32，SM最多放32个block（最大block数限制），总共1024线程，占用率只剩50%。

这说明什么？在128到512这个区间，占用率和block大小基本解耦。真正影响占用率的是寄存器用量和共享内存用量。如果一个kernel每个线程用了超过32个寄存器，SM的64K寄存器文件只能支持少于2048线程，那么就算block=128，也可能只驻留8个block，占用率只有50%。

高占用率的意义在于隐藏访存延迟。GPU访问全局内存的延迟动辄几百个周期，如果SM上没有足够的warp来切换，那就只能干等。所以访存密集型kernel通常需要高占用率来把延迟填满。但占用率不是越高越好。有些kernel对缓存和shared memory的复用要求很高，block小一点反而能增大单个线程的cache命中率。实测下来，80%到100%占用率的性能差别很多时候只有个位数百分比，没必要为了最后一点占用率把block调得面目全非。

3.4 寄存器与共享内存对驻留block数的影响

寄存器是每个SM的硬资源，有限且宝贵。SM的寄存器文件一般是65536个32位寄存器。如果编译器给一个kernel分配了32个寄存器/线程，那SM最多同时容纳2048个线程（65536/32=2048）。如果每个线程用到64个寄存器，SM最多容纳1024个线程。在这种时候，你即使设置了block=256，SM也只能放4个block（而非8个），占用率跌到50%。

怎么查？编译时加寄存器信息输出：

nvcc -arch=sm_80 --ptxas-options=-v vectorAdd.cu

编译日志会打出每个kernel用掉的寄存器数和shared memory字节数。拿到这个数字，再对照SM的资源总量，才能算出真实的驻留block数。

更稳妥的方法是直接调用运行时API：

int maxActiveBlocks = 0; cudaOccupancyMaxActiveBlocksPerMultiprocessor( &maxActiveBlocks, vectorAdd, blockSize, 0); printf("blockSize=%d, maxActiveBlocks per SM=%d\n", blockSize, maxActiveBlocks);

这个API会把寄存器、shared memory、线程数限制全部考虑进去，直接告诉你给定block大小下每个SM最多能放几个block。调试配置的时候非常方便。

共享内存也一样。每个SM的shared memory容量限制（比如48KB到228KB，不同架构差异很大）决定了能驻留多少个block。如果你在kernel里开了比较大的动态shared memory数组，SM能放下的block数量会减少。这也是为什么shared memory和寄存器用量会成为block大小决策的“隐藏尾巴”——你不能光看线程数在不在1024以内，还得看资源挤不挤得下。

4. 从简单到复杂的实操配置示例

4.1 一维向量加法：标准block/grid计算方法

这是CUDA的hello world，但恰恰最能说明配置过程。我要做的是两个长度为n = 1 << 20（1048576）的float数组加法。

先决定block大小：选256。这个数字是32的倍数，在推荐区间内，看起来平平无奇但实测稳定。

再算grid大小：

int blockSize = 256; int gridSize = (n + blockSize - 1) / blockSize;

n正好是1048576，除以256等于4096，gridSize正好是4096，没有余数。如果n不是blockSize的整数倍，这个公式会向上取整，保证覆盖所有数据，代价是最后一个block有一批线程不做有效工作。

kernel长这样：

__global__ void vectorAdd(const float* A, const float* B, float* C, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { C[idx] = A[idx] + B[idx]; } }

如果少了if (idx < n)这个边界判断会怎样？假设n=1000000，block=256，gridSize=3907，最后一个block的线程数只用了3907*256-1000000=192个线程，其余64个线程会去写C[1000000]到C[1000063]这段非法内存。后果可能是数据损坏，也可能是段错误，而且错误报出来的位置往往离实际出错点很远，排错很痛苦。

启动之后必须检查错误。kernel启动是异步的，普通错误不会立刻抛出，所以在启动后紧跟着调用：

CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaDeviceSynchronize());

我习惯把错误检查封装成宏：

#define CUDA_CHECK(call) \ do { \ cudaError_t err = (call); \ if (err != cudaSuccess) { \ printf("CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \ cudaGetErrorString(err)); \ exit(1); \ } \ } while (0)

这样一旦配置有问题，日志能直接定位到文件行号。

4.2 二维矩阵乘法：block维度怎么拆

矩阵乘法是二维block配置的典型场景。假设要算一个M x K的矩阵乘以K x N的矩阵，输出是M x N。最朴素的想法是：每个线程负责计算输出矩阵的一个元素，也就是C[row][col]。

因为row和col是二维的，block用二维就很自然。经典配置是block(16, 16)，总共256个线程。

__global__ void matMul(const float* A, const float* B, float* C, int M, int N, int K) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; if (row < M && col < N) { float sum = 0.0f; for (int k = 0; k < K; ++k) { sum += A[row * K + k] * B[k * N + col]; } C[row * N + col] = sum; } }

启动参数这样算：

dim3 block(16, 16); dim3 grid((N + 15) / 16, (M + 15) / 16); matMul<<<grid, block>>>(dA, dB, dC, M, N, K);

为什么要16x16，而不是32x32或8x8？

32x32就是1024个线程，已经触及block最大线程数。把它设成32x32虽然合法，但一个SM最多只能放2个这样的block，调度灵活性急剧下降。而且矩阵乘法通常要配合shared memory做分块，32x32的tile需要1024个float的shared memory，再叠加寄存器压力，很可能让SM驻留不了几个block，占用率崩掉。

8x8这种小block只有64线程，tile切得太碎，每个线程只负责很小一块数据，但块数特别多，调度开销上升，shared memory的复用率也下降。16x16是几十年来GPU社区验证过的平衡点：块内线程数256，warp映射规整，一个warp覆盖两行数据的访问模式还比较集中，shared memory tile大小也适中。

如果你要更精细地调矩阵乘法，后边肯定要往tile分块、寄存器缓存、向量化这些方向走，但第一步把grid和block配成16x16是稳妥的起点。换别的block尺寸（比如block(32, 8)）也不是不行，只是你需要用profiling数据说话，而不是拍脑袋。

4.3 什么时候才考虑1024线程的大block

既然上限是1024，那什么时候真的用它？

我在实践中发现，大block适合那种“一个block就是一个完整的工作组”的场景。比如某些基于block的计算，每个block里需要共享一份较大的数据，或者需要大量线程做归约（reduction）。block内做归约时，1024个线程分两步就能归并完：第一步每个warp内部归约，第二步32个warp的结果再归约。线程越多，单block能处理的数据越多，跨block协作的开销越小。

还有一种场景是每个线程的工作量极小，比如只是读一个值、做一次比较、写一个结果。这个时候用大block可以让kernel的总线程数少一些，减少调度的指令开销。但这类kernel通常不是性能瓶颈，纠结意义不大。

如果你拿不准，我的建议很简单：默认256或128，完事之后用Nsight出几组对比。CUDA调参没有一劳永逸的答案，但有很宽很稳的默认区。

5. 常见配置陷阱与排查经验

5.1 启动失败先看cudaErrorInvalidValue

kernel启动失败，最常见的错误是cudaErrorInvalidValue。这个错误说明启动参数不合法，通常不外乎这几种情况：

• block线程数超过1024
• blockDim.z超过64
• grid的y或z方向超过65535
• 传了空指针或非法内存地址

排查思路是缩小范围。先把block降到256，grid只保留x方向，跑一遍。如果能跑，说明问题出在维度配置上。然后逐个维度往回去。这种二分法在CUDA排错里特别实用。

注意，kernel启动报错往往是异步的，可能到cudaMemcpy或者cudaDeviceSynchronize才把错误吐出来。所以启动后立刻cudaGetLastError()是很有用的习惯，能避免错报位置误导你。

5.2 性能差的排查思路要系统化

如果你把配置改来改去，性能还是上不去，不要继续盲猜block大小。正确的顺序是：

先看占用率。用cudaOccupancyMaxActiveBlocksPerMultiprocessor或Nsight Compute看SM实际驻留了多少warp。如果占用率偏低，优先查寄存器数和shared memory是不是超了。寄存器超了，代码里加__launch_bounds__(256, 8)这种约束强行限流，或者在编译选项里调maxrregcount。

再看访存模式。直接用Nsight Compute看是否有大量的非合并访问。这往往不是block大小导致的，而是你的数据索引映射错了。比如前面提到的把col和row记反，或者用一维block去处理本来应该二维连续的图像数据。

最后对比几组常规配置。我的套路是用blockSize在64、128、256、512、1024这五档各跑一遍，记录耗时，画个表。绝大多数情况下，128或256会胜出，如果512和1024有明显优势，再去深挖原因。这样至少能从数据上确定方向，而不是瞎猜。

5.3 边界条件漏判和分支分歧

边界条件漏判是最隐蔽的问题之一。调试时数据量小，比如n=1000，整除正好，边界条件永远不触发，看起来一切正常。一到正式数据量，n=1000000，最后一个block越界，程序在release模式下崩溃，措手不及。

所以从一开始就把边界判断写进kernel，不管数据能不能整除。

分支分歧的问题比较微妙。如果你的数据按连续区域划分，比如每个线程处理i和i + stride，那么同一个warp里的线程索引是连续的，边界判断造成的分支分歧只发生在最后那个warp，无伤大雅。如果你把线程ID重排成某种跳跃映射，比如按block的y方向做列访问，那同一个warp里可能大量线程走不同分支，性能损失就大了。

5.4 善用工具定位occupancy问题

最后强烈推荐两个工具，一个是cudaOccupancyMaxActiveBlocksPerMultiprocessor，一个是Nsight Compute。

前者适合在代码里快速验证配置；后者适合启动完整个程序后，看每个kernel的SM利用率、内存吞吐、warp stall原因。我见过太多人调CUDA性能时只盯着grid和block，但真正拖后腿的可能是一个意外的shared memory bank conflict，或者一个没有展开的循环。Nsight Compute能把这些问题列得清清楚楚。

调参时我的个人习惯是：先保证配置合法、数据不越界、没有合并访问问题，然后再用工具看occupancy和stall原因。最后再动手改block大小，每改一次跑一次性能对比。通常跑三轮就能确定一个稳定配置，比纸上谈兵强得多。

说实话，我最开始也被grid、block、thread三个概念绕晕过，还天真地找过所谓“最优配置”。后来调多了才明白，它本质上是一个在硬件限制范围内的tuning问题。每个硬件有它的脾气，每个kernel有它的资源占用特征，没有放之四海皆准的数字。但只要你把warp对齐、边界检查、资源限制这几个底层逻辑想清楚，把128/256/512这几档跑明白，再配合Nsight看两轮数据，配置block就成了一件有章可循的事。