2. CUDA编程回顾：CUDA线程模型、内存层次结构、核函数编写基础
好，咱们正式开始第二讲。说实话，很多做TensorRT部署的同学，一遇到性能瓶颈就慌了，到处调参数。但在我看来，不懂CUDA底层，调参就是瞎蒙。这一节，我带大家快速回顾一下CUDA编程最核心的三个东西：线程模型、内存层次、核函数写法。这些都是后面做联合调试的基石。
2.1 CUDA线程模型：别把线程当"人"看
很多初学者会把GPU线程理解成CPU线程，这是个大坑。我刚开始学的时候也犯过这个错。CPU线程是 heavyweight 的，切换开销大。但GPU线程呢？轻量到几乎可以忽略不计。
CUDA的线程模型是分层的，说白了就是三级结构：
- Thread（线程）：最基础的执行单元
- Block（线程块）：一组线程，可以协作，共享内存
- Grid（网格）：一组线程块，构成整个任务
你想想看，一个Grid里可以有成千上万个Block，一个Block里又可以有几百个Thread。这玩意儿怎么调度？
关键点：GPU的调度单位其实是 Warp（线程束），32个线程一组。同一个Warp里的线程，执行的是同一套指令。如果分支分歧了，那性能就惨了。我在项目中遇到过，一个简单的if-else没处理好，直接让吞吐量掉了30%。
写核函数的时候，我们通常这样获取线程索引：
// 获取全局线程ID
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 如果是二维网格
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;
int global_idx = idy * gridDim.x * blockDim.x + idx;我个人习惯把 blockIdx 和 threadIdx 当成坐标来理解。每个线程都知道自己是谁，该处理哪块数据。这就是所谓的 SIMT（单指令多线程） 模型。
2.2 内存层次结构：速度与容量的博弈
CUDA的内存模型，说白了就是一套金字塔结构。越往上越快，但容量越小。我经常跟团队说：搞懂内存层次，你就掌握了CUDA性能调优的一半。
| 内存类型 | 位置 | 访问速度 | 作用域 | 生命周期 |
|---|---|---|---|---|
| 全局内存 | 显存 | 慢（~200周期） | 所有线程 | 程序运行期间 |
| 共享内存 | 芯片内 | 快（~5周期） | 同一个Block | Block生命周期 |
| 寄存器 | SM内部 | 极快（~1周期） | 单个线程 | 线程生命周期 |
| 常量内存 | 显存（带缓存） | 较快（有缓存） | 所有线程 | 程序运行期间 |
| 纹理内存 | 显存（带缓存） | 较快（空间局部性） | 所有线程 | 程序运行期间 |
嗯，这里要注意：全局内存的访问模式非常关键。如果线程访问的地址是连续的，就能合并访问，一次搞定。如果跳来跳去，那就惨了。我曾经调试过一个矩阵转置的核函数，就是因为没注意合并访问，性能差了10倍。
我的经验：能用共享内存就别用全局内存。共享内存相当于一个手动管理的缓存。把数据从全局内存搬到共享内存，让Block内的线程协作处理，往往能带来数量级的提升。
共享内存的声明方式：
// 静态声明
__shared__ float shared_data[256];
// 动态声明（需要在核函数调用时指定大小）
extern __shared__ float dynamic_shared[];为什么会有动态声明？因为有时候Block大小不确定，你没法写死。我在做动态尺寸的模型部署时，经常用动态共享内存。
2.3 核函数编写基础：从Hello World到实战
核函数就是跑在GPU上的函数。写法上有个硬性要求：必须返回void，前面加 __global__ 修饰符。来，看个最简单的例子：
// 向量加法核函数
__global__ void vecAdd(float* A, float* B, float* C, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
C[i] = A[i] + B[i];
}
}
// 调用方式
int N = 1 << 20;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);这里有个细节：if (i < N) 这个边界检查不能少。为什么？因为 blocksPerGrid * threadsPerBlock 可能大于 N，多余的线程得让它直接返回。我见过有人忘了这个，结果数组越界，查了半天bug。
避坑指南：我曾经在项目里犯过一个低级错误——在核函数里用 printf 调试。嗯，虽然CUDA支持 printf，但输出顺序是乱的，而且会严重影响性能。建议用 cuda-gdb 或者 cuda-memcheck 来调试。
核函数里能用的内置变量：
gridDim：网格的维度（Block的数量）blockIdx：当前Block的索引blockDim：Block的维度（Thread的数量）threadIdx：当前Thread的索引warpSize：Warp的大小（固定32）
再聊一个常见的优化技巧——循环展开。比如做归约求和，如果让每个线程只处理一个元素，那效率太低了。通常的做法是让每个线程处理多个元素：
__global__ void reduce_v1(float* input, float* output, int N) {
extern __shared__ float sdata[];
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x * 2 + threadIdx.x;
// 每个线程加载两个元素
sdata[tid] = (i < N) ? input[i] : 0;
sdata[tid] += (i + blockDim.x < N) ? input[i + blockDim.x] : 0;
__syncthreads();
// 归约操作
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
if (tid == 0) output[blockIdx.x] = sdata[0];
}你看，这里每个线程处理了两个元素，减少了线程总数，提高了利用率。而且用了 __syncthreads() 做同步，保证共享内存里的数据是完整的。
说到同步，__syncthreads() 是个好东西，但别滥用。它会让Block内所有线程都等到同一个点。如果有些线程提前返回了，那就会死锁。我建议只在共享内存读写之间使用。
2.4 小结：打好基础再谈优化
这一节的内容，说白了就是CUDA编程的"内功心法"。线程模型让你知道任务怎么分，内存模型让你知道数据怎么放，核函数编写让你知道代码怎么写。这三样东西，在后面的TensorRT调试中会反复用到。
我个人建议，如果你刚开始接触CUDA，先别急着调优。把上面这些基础写熟练了，能随手写出正确的核函数，再谈性能优化。下一节，我们会讲CUDA的流和事件，那是做异步并发的基础。
嗯，今天就到这儿。有问题随时交流。