3. CUDA编程模型

好，咱们今天聊聊CUDA编程模型里最核心的东西——Grid、Block、Thread这三层结构。说实话，我刚开始接触CUDA时，也被这玩意儿绕得有点晕。但后来在项目中调优调多了，慢慢就摸清了门道。

说白了，CUDA的编程模型就是一套“怎么把任务拆开，再扔给GPU去并行执行”的规则。你想想看，GPU里有几千个核心，你总得告诉它每个核心该干嘛吧？这三层结构就是干这个的。

3.1 三层结构：从宏观到微观

咱们先看一个最直观的对应关系：

• Grid（网格）：对应一个完整的计算任务。比如你要处理一张1024x1024的图片，这就是一个Grid。
• Block（线程块）：Grid被拆成若干Block。每个Block负责图片中的一个区域，比如64x64像素。
• Thread（线程）：Block里包含多个Thread。每个Thread处理一个像素点。

嗯，这里要注意：Thread才是真正干活的最小单位。Grid和Block更多是逻辑上的组织方式。

我在项目中遇到过一个问题：有同事把Block设得特别大，结果性能反而下降了。为什么？因为硬件对Block的大小有限制。咱们后面会细说。

3.2 硬件映射：逻辑结构如何落地

好，现在咱们看看这三层结构是怎么映射到GPU硬件的。我画了一张图，帮你理清关系：

这张图你看懂了吗？我解释一下：

• Grid 是逻辑概念，对应一次kernel启动
• Block 会被调度到 SM（流多处理器） 上执行
• Thread 最终在 CUDA Core 上运行

这里有个关键点：一个Block只能在一个SM上执行，但一个SM可以同时运行多个Block。我当年调优时，就靠这个特性把性能翻了一倍。

3.3 线程索引：怎么找到自己的数据

每个Thread都需要知道自己该处理哪部分数据。CUDA提供了内置变量：

// 一个典型的CUDA kernel
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    // 计算全局线程ID
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

你看，blockIdx.x 是当前Block的索引，threadIdx.x 是Thread在Block内的索引，blockDim.x 是Block的大小。三者一组合，就得到了全局唯一的线程ID。

3.4 硬件限制：别踩这些坑

嗯，这里要重点说说硬件限制。每个SM的资源是有限的：

资源	典型限制	影响
最大Thread数/Block	1024	Block不能太大
最大Block数/SM	16-32	Block太多会排队
共享内存大小	48KB-96KB	共享内存用多了，Block数就少了
寄存器数量	65536/SM	每个Thread用的寄存器多了，能跑的Thread就少了

所以，调优时你要平衡这几个因素：

• Block大小：建议是32的倍数，256或512比较稳妥
• Grid大小：至少是SM数量的2倍，保证负载均衡
• 资源占用：别让每个Thread用太多寄存器和共享内存

3.5 实际案例：矩阵乘法

咱们看个实际例子。假设你要计算两个1024x1024矩阵的乘法：

// 定义Block和Grid大小
dim3 blockDim(16, 16);  // 每个Block 16x16 = 256个线程
dim3 gridDim(64, 64);   // 64x64 = 4096个Block

// 启动kernel
matMulKernel<<<gridDim, blockDim>>>(d_A, d_B, d_C, 1024);

这里每个Block处理输出矩阵中的一个16x16小块。为什么选16x16？因为：

• 256个线程，是32的整数倍，warp利用率高
• 每个线程只需要少量寄存器，SM能同时跑更多Block
• 共享内存刚好够用，不会成为瓶颈

我个人习惯是，先按这个配置跑一遍，然后用NVIDIA的Nsight工具看看occupancy（占用率）。如果occupancy低于50%，就调小Block或减少资源使用。

3.6 小结

好，咱们总结一下今天的内容：

• Grid、Block、Thread 是CUDA的三层逻辑结构，用来组织并行任务
• Block映射到SM，Thread映射到CUDA Core，这是硬件层面的对应关系
• 线程索引 通过 blockIdx、threadIdx、blockDim 计算
• 硬件限制 决定了Block大小和资源使用的上限

记住一句话：逻辑上怎么拆，硬件上怎么跑。搞懂了这三层结构，你就掌握了CUDA编程的骨架。剩下的，就是往里面填肉了。