GPU编程入门：CUDA基础、GPU架构、核函数、内存层次结构
说实话，我第一次接触GPU编程的时候，心里是有点发怵的。那时候我刚从嵌入式C转过来，满脑子都是CPU那套“单核串行”的思维。结果一上来就面对几千个线程同时跑，感觉就像从开手动挡小车，突然跳进了波音747的驾驶舱。但后来我发现，只要你理解了GPU的“脾气”，它其实比CPU更直来直去。
1. GPU架构：为什么它和CPU不一样？
CPU和GPU的设计哲学，说白了就是“精英小队” vs “人海战术”。CPU有强大的控制单元和巨大的缓存，擅长处理复杂的逻辑分支。而GPU呢？它把大部分晶体管都用来做计算单元（ALU），控制单元和缓存都精简到极致。
我习惯把GPU看作一个“包工头”带着几千个“民工”。包工头（控制单元）只负责分配任务，民工们（CUDA核心）只管埋头干活。每个民工虽然单兵作战能力一般，但架不住人多啊！
核心差异：
- CPU： 延迟优化 —— 想尽办法让单个任务跑得快。
- GPU： 吞吐量优化 —— 不在乎单个任务慢一点，但整体要完成海量任务。
你想想看，如果你要算100万个数的加法，CPU会一个一个算，但GPU会分成1000组，每组1000个数同时算。这就是并行计算的魅力。
2. CUDA基础：从Hello World开始
CUDA是NVIDIA推出的并行计算平台。你不需要把它想得太复杂，它就是在C语言基础上加了一些扩展关键字。比如 __global__ 表示这个函数在GPU上运行，__device__ 表示在GPU上调用。
来看一个最简单的核函数：
// 核函数定义：每个线程执行一次
__global__ void vec_add(float *a, float *b, float *c, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
int main() {
int n = 1024;
float *d_a, *d_b, *d_c;
// 在GPU上分配内存
cudaMalloc(&d_a, n * sizeof(float));
cudaMalloc(&d_b, n * sizeof(float));
cudaMalloc(&d_c, n * sizeof(float));
// 启动核函数：1个block，1024个线程
vec_add<<<1, 1024>>>(d_a, d_b, d_c, n);
cudaDeviceSynchronize();
// 记得释放内存
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
return 0;
}个人习惯： 我每次写核函数，都会在开头加一个边界检查 if (idx < n)。因为线程总数不一定能整除数据量，多出来的线程必须“空转”返回。曾经我漏掉这个检查，结果数组越界，查了整整一下午。
3. 核函数：线程是怎么组织的？
GPU的线程组织分三层：网格（Grid）→ 线程块（Block）→ 线程（Thread）。你启动核函数时写的 <<<grid, block>>> 就是在定义这两层。
- Grid： 包含所有Block。可以是一维、二维、三维。
- Block： 包含一组Thread。同一个Block内的线程可以协作（共享内存、同步）。
- Thread： 最小编程单元，每个线程执行核函数的一份拷贝。
为什么这样设计？说白了是为了适应不同维度的数据。比如处理一张图片，你可以用二维Grid（对应图像宽高），每个Block处理一个16x16的像素块。
内置变量：
threadIdx.x：线程在Block内的索引blockIdx.x：Block在Grid内的索引blockDim.x：Block的维度大小gridDim.x：Grid的维度大小
嗯，这里要注意：一个Block内的线程数上限通常是1024（取决于GPU架构）。我刚开始写的时候，总想一个Block塞2048个线程，结果程序直接崩了。后来才养成习惯，先查一下 cudaDeviceGetAttribute 获取设备参数。
4. 内存层次结构：谁快谁慢，心里要有数
GPU的内存体系，是性能调优的关键。我见过太多人把所有数据都放在全局内存里，结果性能还不如CPU。说白了，你得知道每种内存的“脾气”。
| 内存类型 | 位置 | 访问速度 | 作用域 | 生命周期 |
|---|---|---|---|---|
| 全局内存 | 板载显存 | 慢（~200周期） | 所有线程 | 程序运行期 |
| 共享内存 | 芯片内（SM） | 快（~5周期） | 同一个Block | Block生命周期 |
| 寄存器 | 芯片内（每个线程） | 极快（1周期） | 单个线程 | 线程生命周期 |
| 常量内存 | 板载显存（带缓存） | 较快（有缓存） | 所有线程（只读） | 程序运行期 |
| 纹理内存 | 板载显存（带缓存） | 较快（空间局部性） | 所有线程（只读） | 程序运行期 |
避坑指南： 我曾经在项目里把一个大数组放在全局内存，每个线程都频繁读写，结果带宽被占满，性能惨不忍睹。后来我把频繁访问的数据搬到了共享内存，速度提升了8倍。记住：尽量减少全局内存访问，能用共享内存就别用全局内存。
5. 共享内存与同步：Block内的“小团队协作”
共享内存是GPU编程的“杀手锏”。它允许同一个Block内的线程互相通信。但有个前提：你必须手动管理同步，不然会出现数据竞争。
来看一个经典的归约求和例子：
__global__ void reduce_sum(float *input, float *output, int n) {
__shared__ float sdata[256]; // 共享内存，大小等于Block线程数
int idx = threadIdx.x + blockIdx.x * blockDim.x;
int tid = threadIdx.x;
// 加载数据到共享内存
sdata[tid] = (idx < n) ? input[idx] : 0.0f;
__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() 时，都要想清楚“所有线程是否都会到达这里？”如果有些线程提前return了，那同步就会死锁。我早期就踩过这个坑，调试时卡死，最后发现是条件分支导致部分线程没执行同步。
6. 内存合并访问：让带宽跑满的秘诀
全局内存虽然慢，但如果你能让访问模式“合并”，就能最大化利用带宽。所谓合并访问，就是相邻的线程访问相邻的地址。
举个例子：
- 好：
c[idx] = a[idx] + b[idx];—— 线程0访问地址0，线程1访问地址1，连续。 - 坏：
c[threadIdx.x * 64] = ...—— 线程0访问地址0，线程1访问地址64，跨度大，无法合并。
我在项目中遇到过一位同事，他把数据结构设计成 struct of arrays 而不是 array of structs，结果性能直接翻倍。说白了，就是让数据在内存中按线程访问顺序紧密排列。
总结一下GPU编程的“心法”：
- 先想清楚数据怎么分块（Block划分）。
- 尽量用共享内存做中间计算。
- 保证全局内存访问是合并的。
- 减少线程分支分歧（warp divergence）。
- 用
cuda-memcheck检查越界。
好了，这一章的内容就到这里。GPU编程其实没那么神秘，你只要把“并行”两个字刻在脑子里，多写多练，很快就能上手。记住，遇到性能瓶颈时，先看看内存访问模式，再查查共享内存使用，往往能解决80%的问题。
公众号:蓝海资料掘金营，微信deep3321