24. GPU编程基础：CUDA入门、数据传输优化、核函数优化
说到GPU编程，很多人第一反应就是“难”。其实不然。我刚开始接触CUDA时，也觉得这东西玄乎——CPU跑得好好的，干嘛非要去折腾显卡？直到有一次，我接手了一个图像处理项目，单张图片要处理3秒钟，客户直接摔了鼠标。嗯，从那以后，我就老老实实学起了CUDA。
说白了，GPU就是一群“头脑简单、四肢发达”的工人。CPU是几个全能博士，什么都能干，但数量少。GPU是成千上万个只会做加减乘除的壮汉，干不了复杂活，但人多力量大。你想想看，处理一张图片的每个像素，不就是让每个壮汉算一个像素值吗？这活他们最擅长。
24.1 CUDA编程模型入门
先说说CUDA的基本概念。我个人习惯把CUDA程序分成两部分：Host端（CPU负责）和Device端（GPU负责）。Host端管逻辑控制、数据准备，Device端管大规模并行计算。
一个典型的CUDA程序长这样：
// 1. 分配GPU内存
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));
// 2. 把数据从CPU拷贝到GPU
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
// 3. 调用核函数（在GPU上执行）
myKernel<<<gridDim, blockDim>>>(d_data, N);
// 4. 把结果拷回CPU
cudaMemcpy(h_result, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
// 5. 释放GPU内存
cudaFree(d_data);这里有个关键点：<<<gridDim, blockDim>>> 这个尖括号里的东西，决定了有多少线程在干活。gridDim 是线程块网格的维度，blockDim 是每个线程块里的线程数。我刚开始写的时候，经常搞混这两个参数，结果要么线程不够用，要么线程太多把GPU撑爆了。
核心概念速查表：
| 概念 | 说明 | 我的经验 |
|---|---|---|
| Thread | 最小的执行单元 | 每个线程执行同一个函数，但处理不同数据 |
| Block | 一组线程，共享共享内存 | block内线程可以协作，但别超过1024个 |
| Grid | 一组Block，构成整个任务 | Grid大小决定了总线程数 |
| Kernel | 在GPU上运行的函数 | 用 __global__ 修饰，返回void |
24.2 数据传输优化：别让带宽成为瓶颈
我在项目中遇到过最坑的事：辛辛苦苦写好了核函数，跑起来发现比CPU还慢。查了半天，发现90%的时间都花在数据传输上。你想想看，PCIe 3.0 x16的理论带宽也就16GB/s，而GPU显存带宽动辄几百GB/s。数据搬运的速度，根本喂不饱GPU这个“大胃王”。
怎么优化？我总结了几个实战技巧：
- 减少传输次数：能一次传完的数据，别分十次传。我见过有人循环里每次传一个float，简直是在折磨PCIe总线。
- 使用异步传输：
cudaMemcpyAsync配合流（Stream），可以让数据传输和计算重叠。说白了就是一边拷数据，一边算数据，别闲着。 - 固定内存（Pinned Memory）：用
cudaHostAlloc分配主机内存，能提高传输带宽。我实测过，能快2-3倍。 - 零拷贝（Zero-Copy）：对于只读数据，可以直接让GPU访问主机内存，省掉拷贝这一步。但要注意，这招只适合数据量小、访问频率低的情况。
避坑指南：我曾经在项目里用了 cudaMallocHost 分配了超大块固定内存，结果系统内存被锁死，其他进程直接卡死。固定内存虽然快，但别贪多，够用就行。
24.3 核函数优化：榨干GPU的每一分性能
核函数是GPU编程的灵魂。写得好，性能飞起；写得烂，GPU在摸鱼。我见过太多人把CPU的写法直接搬到GPU上，结果惨不忍睹。
先看一个简单的向量加法核函数：
__global__ void vecAdd(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}这个写法没问题，但性能一般。怎么优化？
24.3.1 内存访问合并
GPU从显存读数据时，一次会读连续的一大块。如果线程访问的内存地址是连续的，就能合并成一次读取，效率极高。反之，如果跳着访问，每次只读一点点，带宽就浪费了。
我建议：让相邻线程访问相邻地址。比如线程0访问a[0]，线程1访问a[1]，这样就是合并访问。如果线程0访问a[0]，线程1访问a[1000]，那就是灾难。
24.3.2 共享内存的使用
共享内存是GPU上的一块“高速缓存”，比全局内存快几十倍。但它的容量很小，一个Block通常只有几十KB。怎么用？
- 把频繁访问的全局数据先搬到共享内存里
- 在Block内部做数据复用，减少全局内存访问
- 注意Bank Conflict：多个线程同时访问共享内存的同一个Bank，会导致串行化
警告：共享内存用多了，每个Block能启动的线程数就会减少。我见过有人为了用共享内存，把Block线程数从512降到128，结果并行度下降，性能反而更差。这玩意是把双刃剑。
24.3.3 线程束发散
GPU执行指令时，是以32个线程为一组（称为Warp）来执行的。如果Warp里的线程走了不同的分支（比如if-else），那这些分支就得串行执行，性能直接腰斩。
我有个习惯：尽量让Warp内的线程走同一个分支。如果实在避免不了，就把分支条件设计成基于 threadIdx.x / 32 这种粒度，让整个Warp保持一致。
24.4 知识体系总览
下面这张图，是我自己总结的CUDA优化知识体系。每次做性能调优时，我都会对照着检查一遍：
24.5 实战经验总结
最后，分享几个我踩过的坑，希望能帮你少走弯路：
- 别迷信“线程越多越好”：GPU的并行度受限于SM数量和寄存器资源。我曾经把线程数设到几百万，结果每个线程分到的寄存器太少，性能反而下降。
- 用好
cuda-gdb和nvprof：我习惯先用nvprof看瓶颈在哪，是计算受限还是带宽受限，再针对性优化。别瞎猜，用数据说话。 - 注意浮点精度：GPU的float运算和CPU不完全一样。我在一个科学计算项目里，因为精度问题导致结果偏差，排查了整整两天。后来改用double，虽然慢了点，但结果对了。
- 别忽略错误检查：每个CUDA API调用后，记得检查返回值。我见过太多人写
cudaMalloc不检查，结果内存不够直接崩了，还找不到原因。
一个小技巧：写CUDA代码时，我习惯在核函数开头加一句 if (threadIdx.x == 0 && blockIdx.x == 0) printf(...)，用来打印调试信息。虽然慢，但关键时刻能救命。
GPU编程说难不难，说简单也不简单。核心就三点：数据搬得快、线程用得好、架构摸得清。把这三点吃透了，大部分性能问题都能迎刃而解。嗯，今天就聊到这，下次有机会再聊聊更高级的优化技巧。
公众号:蓝海资料掘金营，微信deep3321