核函数启动配置：Grid、Block、Thread的维度设计原则

各位同学，今天我们来聊聊核函数启动配置。说白了，就是怎么安排你的线程大军去干活。

我刚开始做CUDA优化时，总觉得这玩意儿随便设设就行。后来被现实狠狠教育了一回——同样的算法，启动配置不同，性能能差出5倍。从那以后，我再也不敢小看这个环节了。

一、线程层次结构：从大到小看

GPU的线程组织分三层：Grid、Block、Thread。你可以想象成一个公司：

• Grid：整个公司，对应一个核函数调用
• Block：部门，对应一组协作的线程
• Thread：员工，真正干活的单元

启动核函数时，我们这样写：

// 定义Grid和Block的维度
dim3 gridDim(16, 16);    // 16x16 = 256个Block
dim3 blockDim(32, 8);    // 32x8 = 256个线程/Block

// 启动核函数
myKernel<<<gridDim, blockDim>>>(d_data, N);

嗯，这里要注意：dim3是三维的，但如果你只传两个参数，第三个默认为1。我习惯把不用的维度设为1，代码更清晰。

二、维度设计原则：别瞎设，有讲究

我个人总结了三条核心原则，你记好了：

三、Occupancy计算：到底能跑多快？

Occupancy，直译是「占用率」。说白了就是：你的SM上同时跑了多少个warp，占理论最大值的比例。

计算公式长这样：

Occupancy = 实际活跃warp数 / 最大支持warp数

举个例子，假设你的GPU每个SM最多支持64个warp（2048线程）：

• 你设了blockDim=256，每个Block有8个warp
• 如果每个SM能放8个Block，那就是64个warp
• Occupancy = 64/64 = 100%

听起来100%最好？不一定。我踩过这个坑——为了追求高Occupancy，把Block设得很小，结果每个线程能用的寄存器少了，反而导致寄存器溢出到全局内存，性能暴跌。

四、Occupancy优化实战：三步走

我一般按这个流程来调：

1. 查硬件参数：用cudaGetDeviceProperties获取SM的warp数、寄存器数、共享内存大小
2. 算理论Occupancy：用NVIDIA的Occupancy Calculator，或者自己算
3. 实际跑分：调Block大小，看实际吞吐量

这里有个小工具，我经常用：

// 获取设备属性
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);

printf("每个SM最大线程数: %d\n", prop.maxThreadsPerMultiProcessor);
printf("每个SM最大Block数: %d\n", prop.maxBlocksPerMultiProcessor);
printf("每个Block最大线程数: %d\n", prop.maxThreadsPerBlock);
printf("寄存器数量: %d\n", prop.regsPerBlock);

拿到这些参数后，你就可以算Occupancy了。举个例子：

Block大小	每个SM的Block数	活跃warp数	Occupancy	实际性能
128	16	64	100%	慢（寄存器不足）
256	8	64	100%	中等
512	4	64	100%	快（寄存器充足）

你看，同样100% Occupancy，性能却不同。关键在寄存器压力。

五、知识体系总览

下面这张图，是我自己画的，帮你理清本章的核心逻辑：

六、实战建议：我踩过的坑

最后，分享几个实战经验：

嗯，今天就讲到这里。记住：启动配置不是玄学，是科学。多算、多测、多调，你也能成为优化高手。