4、CUDA基础回顾：编程模型、内存层次、线程块配置与性能工具
好，咱们进入第四章。这一章我打算带大家快速过一遍CUDA的核心基础。别觉得基础就简单，我在边缘设备上踩过的坑，十有八九都是基础没打牢。
说白了，CUDA编程就是让GPU干活。但GPU怎么干活，你得先懂它的规矩。
4.1 CUDA编程模型：Host与Device的协作
CUDA的编程模型，核心就两个角色：Host（CPU）和Device（GPU）。Host负责逻辑控制、数据准备，Device负责大规模并行计算。
我习惯把Host比作项目经理，Device比作施工队。项目经理把图纸（数据）和任务（kernel）交代清楚，施工队就闷头干。干完了，项目经理去验收（数据回传）。
典型的CUDA程序流程是这样的：
- 分配Host内存，初始化数据
- 分配Device内存
- 把数据从Host拷贝到Device
- 调用kernel函数，在GPU上执行
- 把结果从Device拷回Host
- 释放内存
嗯，这里有个关键点：kernel是异步执行的。也就是说，你调用完kernel，CPU不会傻等着，它会继续往下走。我曾经在项目里犯过这个错——kernel还没算完，CPU就去读结果了，读回来的全是垃圾数据。后来加了cudaDeviceSynchronize()才搞定。
核心概念：CUDA编程本质上是异构编程。你要时刻清楚，哪段代码跑在CPU上，哪段跑在GPU上。数据在哪，计算就在哪。
4.2 内存层次结构：别让数据成为瓶颈
GPU的内存层次，是性能优化的重中之重。你想想看，计算单元再快，数据喂不进去也是白搭。
CUDA的内存模型大致分这几层：
| 内存类型 | 位置 | 访问速度 | 作用域 | 生命周期 |
|---|---|---|---|---|
| 全局内存 | 板载显存 | 慢（几百周期） | 所有线程 | 程序全程 |
| 共享内存 | 芯片内 | 快（几个周期） | 同一个Block | Block生命周期 |
| 寄存器 | SM内部 | 极快（1周期） | 单个线程 | 线程生命周期 |
| 常量内存 | 板载显存（有缓存） | 较快 | 所有线程（只读） | 程序全程 |
| 纹理内存 | 板载显存（有缓存） | 较快 | 所有线程（只读） | 程序全程 |
我个人最常用的优化手段，就是把频繁访问的全局数据搬到共享内存里。举个例子，做矩阵乘法时，如果每个线程都去全局内存读数据，那带宽很快就被打满了。但如果你让一个Block的线程先把数据拉到共享内存里，再各自计算，速度能快好几倍。
避坑指南：我曾经在边缘设备上做图像处理，没注意共享内存的大小限制。一个Block开了太多共享内存，结果每个SM能同时跑的Block数量骤降，并行度反而下降了。记住：共享内存不是越大越好，要平衡。
4.3 线程块配置：找到那个黄金比例
线程块怎么配？这是每个CUDA新手都会问的问题。我的回答是：没有万能公式，但有经验法则。
先看几个关键参数：
- 线程块大小（blockDim）：通常是32的倍数，比如128、256、512。为什么？因为warp是32个线程一组，32的倍数能充分利用。
- 网格大小（gridDim）：根据数据量来定。比如你要处理100万个元素，每个线程处理1个，那网格大小就是100万 / 线程块大小。
- 每个SM的最大线程数：不同架构不一样。Maxwell是2048，Pascal是2048，Volta是1024。你得查手册。
我一般这样配：先定线程块大小为256。为什么是256？因为256个线程需要的寄存器数量适中，共享内存也适中，不容易撞到资源上限。然后根据数据量算网格大小。
举个例子：
// 假设有N个元素，每个线程处理1个
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
myKernel<<<gridSize, blockSize>>>(d_data, N);
但注意，这只是起点。实际项目中，我会用性能分析工具跑一遍，看看occupancy（占用率）是多少。如果occupancy太低，就调小线程块或者减少每个线程的寄存器使用量。
警告：线程块大小不要小于32。否则一个warp里只有部分线程在工作，剩下的都在摸鱼，浪费计算资源。我见过有人设blockSize=8，结果性能惨不忍睹。
4.4 流与事件：让GPU流水线跑起来
流（Stream）是CUDA里实现并发的手段。说白了，就是让GPU同时干几件事。
默认情况下，所有操作都在默认流里，是串行的。但你可以创建多个流，让数据传输和计算重叠起来。
我常用的模式是这样的：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在stream1里做数据传输
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
// 在stream2里做计算
kernel1<<<grid, block, 0, stream2>>>(d_data2);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
事件（Event）则用来做同步和计时。比如你想知道某个kernel跑了多久：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
myKernel<<<grid, block>>>(d_data);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("Kernel用时: %f ms\n", milliseconds);
嗯，这里要注意：事件记录也是异步的。你得用cudaEventSynchronize等它完成，才能读时间。我刚开始用的时候忘了同步，读出来的时间全是0，还以为是GPU太快了...
4.5 性能分析工具：别靠猜，靠数据
优化性能最忌讳的就是「我觉得」。你得用工具说话。
我常用的工具有这几个：
- nvidia-smi：快速查看GPU利用率、显存占用、温度。适合现场排查。
- nvprof：老牌性能分析工具，能告诉你每个kernel的耗时、数据传输量、occupancy等。不过NVIDIA已经停止更新了，推荐用Nsight Systems。
- Nsight Systems：新一代工具，界面更友好。能展示时间轴，让你看到CPU和GPU的活动是怎么重叠的。
- Nsight Compute：专门分析kernel性能的。能告诉你瓶颈在计算还是访存，还能给出优化建议。
我个人的工作流是这样的：先用Nsight Systems看整体，找到最耗时的kernel。然后用Nsight Compute深入分析那个kernel，看看是计算密集还是访存密集，再针对性优化。
一个小技巧：在边缘设备上，nvidia-smi可能没有图形界面。你可以用watch -n 1 nvidia-smi每秒刷新一次，实时监控GPU状态。我在调试一个嵌入式设备时，就是靠这个命令发现显存泄漏的。
好了，这一章的内容就到这里。CUDA基础虽然看起来简单，但每个点都值得深挖。下一章我们会把这些知识用到边缘AI设备的实际部署中，到时候你会发现，今天讲的内存层次和线程块配置，都是决定模型能不能跑起来的关键。