一、CUDA编程模型回顾：从GPU架构到程序执行
各位同学好，我是你们的老朋友。今天咱们来聊聊CUDA编程模型的基础。说实话，这部分内容看似简单，但很多人在实际项目中栽跟头，往往就是基础没打牢。我个人习惯把这一章叫做「地基」，地基不稳，后面写再多代码也是空中楼阁。
1.1 GPU架构概览：它到底强在哪？
先问大家一个问题：为什么GPU能比CPU快那么多？
你想想看，CPU就像一位全能博士，什么都会，但一次只能处理几件事。GPU呢？它更像是一群小学生，每个人只会简单的加减法，但胜在人多——成千上万个计算单元同时干活。
具体来说，NVIDIA的GPU架构有几个关键点：
- SM（Streaming Multiprocessor）：核心计算单元，每个SM包含多个CUDA Core
- Warp：32个线程为一组，这是GPU调度和执行的最小单位
- 大量寄存器：每个SM有数万个寄存器，但注意这是共享资源
核心认知：GPU的强项是「大量并行、简单计算」。如果你让GPU做复杂的逻辑分支判断，那效率反而会下降。我在项目中遇到过有人拿GPU做字符串解析，结果性能还不如CPU——这就是典型的用错场景。
下面这张图是我自己画的，帮你理清GPU的硬件层次与软件层次的对应关系：
1.2 线程层次结构：Grid-Block-Thread
CUDA把线程组织成三层结构，这个设计非常巧妙。我刚开始学的时候觉得这有点多余，后来才发现——没有这个层次，你根本没法高效利用GPU。
- Thread（线程）：最基础的计算单元，每个线程执行相同的kernel函数
- Block（线程块）：一组线程的集合，共享Shared Memory，可以同步
- Grid（网格）：所有Block的集合，一个kernel启动对应一个Grid
我的经验：Block大小一般设为128或256比较稳妥。太小了SM利用率低，太大了寄存器不够用。我曾经在项目里设了1024的Block大小，结果每个线程能用的寄存器少得可怜，性能反而下降——这就是典型的「贪多嚼不烂」。
来看一个简单的代码示例，感受一下线程索引的计算：
// 一个简单的向量加法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在Grid中的索引，threadIdx.x 是Thread在Block中的索引，blockDim.x 是Block的维度大小。三者配合，就能唯一定位一个线程。
1.3 内存层次结构：用好每一级存储
GPU的内存层次，说白了就是「速度越快的越少，容量越大的越慢」。这个道理和CPU的缓存类似，但GPU的层次更分明。
| 内存类型 | 作用域 | 延迟（大约） | 容量 | 典型用途 |
|---|---|---|---|---|
| Register | 单个线程 | ~1 cycle | 每个SM约64K~256K | 局部变量、中间结果 |
| Shared Memory | 单个Block | ~5 cycles | 每个SM 48KB~164KB | Block内数据共享、归约 |
| Global Memory | 所有线程 | ~400 cycles | 几GB~几十GB | 主数据存储 |
| Local Memory | 单个线程 | ~400 cycles | 同Global | 寄存器溢出时使用 |
避坑指南：我曾经在优化一个矩阵乘法的kernel时，发现性能始终上不去。查了半天，原来是编译器把一些变量放到了Local Memory里——因为寄存器不够用了。Local Memory虽然对线程是私有的，但它实际上存在Global Memory里，延迟高得吓人。解决办法？减少每个线程的变量数量，或者调整Block大小。
1.4 CUDA程序编译与执行流程
CUDA程序的编译流程，和普通C程序不太一样。它需要经过两阶段编译：
- 设备端编译：.cu文件中的kernel代码被编译成PTX（中间表示）
- 主机端编译：主机代码被C/C++编译器处理
- 链接：最终生成可执行文件，运行时由驱动将PTX编译成机器码
执行流程就更直观了：
// 典型执行流程
1. 分配GPU内存：cudaMalloc()
2. 数据从CPU拷贝到GPU：cudaMemcpy(..., cudaMemcpyHostToDevice)
3. 启动kernel：kernelName<<<grid, block>>>(args)
4. 数据从GPU拷贝回CPU：cudaMemcpy(..., cudaMemcpyDeviceToHost)
5. 释放GPU内存：cudaFree()关键点：kernel的启动是异步的！也就是说，<<<>>> 调用后CPU会立即返回，不会等GPU算完。如果你紧接着就要用结果，必须加 cudaDeviceSynchronize() 同步。我见过太多新手因为这个bug，拿到的数据全是错的。
嗯，到这里第一章的核心内容就差不多了。记住：GPU架构决定了你怎么写代码，内存层次决定了你的性能上限，编译流程决定了你调试的方向。把这些基础吃透了，后面咱们讲优化技巧时，你才能理解「为什么这么改能快」。
想要更多资料，联系可可