思维转变第一步:从“单线程”到“多线程”的认知跃迁
各位同学,欢迎来到第二讲。
上一章我们聊了CPU和GPU的根本差异。今天,我们要跨出最关键的一步——理解SIMT执行模型。说白了,就是搞清楚GPU到底是怎么同时干那么多活的。
我刚开始接触GPU编程时,犯过一个很蠢的错误。我写了一个循环,心想“CPU上跑得挺快,GPU肯定更快”。结果呢?慢得离谱。后来我才明白——思维方式没转过来。
1. 你熟悉的CPU多核:几个大力士
先说说CPU的多核。你想想看,一个4核CPU,就像4个大力士。每个大力士都能独立干活,各干各的。
- 任务并行:核0跑浏览器,核1跑编译器,核2跑视频播放器
- 每个核独立:有自己的缓存、指令流水线、乱序执行单元
- 核数少:桌面CPU通常4-16核,服务器也就几十核
我做过一个项目,要在CPU上并行处理图像。开了4个线程,每个线程处理一块。嗯,效果还行。但当我试图开到16个线程时,性能反而下降了。为什么?因为CPU核就那么几个,线程切换是有开销的。
CPU多核的本质:每个核都是“全能选手”,适合处理复杂、分支多、依赖强的任务。但核数有限,不适合大规模数据并行。
2. GPU的众核:一支蚂蚁军团
GPU就不一样了。它不搞几个大力士,而是搞一支蚂蚁军团——成千上万个轻量级核心。
这些核心有个特点:它们不能独立行动。它们必须听同一个指挥,执行同一条指令。这就是SIMT的核心思想。
| 特性 | CPU多核 | GPU众核 |
|---|---|---|
| 核心数量 | 4-16(少而强) | 数千(多而弱) |
| 每个核心能力 | 强(乱序、分支预测、大缓存) | 弱(顺序执行、小缓存) |
| 并行方式 | 任务并行(各干各的) | 数据并行(一起干同一件事) |
| 适用场景 | 复杂逻辑、串行依赖 | 大规模数据、简单计算 |
我记得第一次看GPU架构图时,密密麻麻的SM(流多处理器)让我头皮发麻。但后来我明白了——每个SM里有一堆CUDA核心,它们共享指令单元。也就是说,32个核心同时执行同一条指令,只是操作的数据不同。
3. SIMT:单指令多线程,到底怎么理解?
SIMT,全称Single Instruction Multiple Threads。名字很绕,但本质很简单:
一条指令,多个线程同时执行,每个线程处理不同的数据。
举个例子。你要给100万个像素做颜色转换。在CPU上,你可能会写一个for循环:
for (int i = 0; i < 1000000; i++) {
pixels[i] = convert(pixels[i]);
}
在GPU上,你写的是这样的:
__global__ void convert_kernel(Pixel* pixels) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
pixels[idx] = convert(pixels[idx]);
}
你看,每个线程只处理一个像素。但GPU会同时启动成千上万个线程,每个线程执行相同的convert函数,只是操作的pixels[idx]不同。
个人经验:我刚开始写CUDA时,总想把循环搬进去。后来发现,在GPU里,循环被“展开”成了线程。你不需要写for循环,你只需要告诉GPU“有多少个线程”,它自己会并行执行。
4. Warp:GPU执行的基本单位
这里有个关键概念——Warp。
GPU并不是一个线程一个线程地调度。它把32个线程打包成一个Warp,一起执行。这32个线程必须执行同一条指令。
为什么会这样?因为硬件设计如此。一个Warp里的线程共享指令单元,节省了芯片面积,提高了效率。
但这也带来了一个问题:分支发散。
我曾经踩过的坑:我在一个Warp里写了if-else分支。结果发现,一半线程走if,一半线程走else。GPU没办法同时执行两条不同指令,只好先执行if,再执行else。性能直接砍半。
所以,写GPU代码时,要尽量避免Warp内的分支。如果一定要分支,尽量让同一个Warp里的线程走同一条路径。
5. CPU vs GPU:思维对比
我们来做个对比,帮你快速切换思维:
- CPU思维:我要优化循环,减少迭代次数,利用缓存局部性
- GPU思维:我要把循环展开成线程,让每个线程处理一小块数据
再举个例子。你要做矩阵乘法。
CPU上,你会写三层嵌套循环,然后考虑循环展开、分块、向量化。GPU上,你会把每个输出元素分配给一个线程,然后考虑怎么让线程高效地读取数据。
说白了,CPU是“少而精”,GPU是“多而糙”。CPU适合处理复杂的控制流,GPU适合处理简单的数据并行。
6. 避坑指南:从CPU迁移到GPU的常见误区
我见过太多人犯这些错误,包括我自己:
- 误区一:线程越多越好
不是的。线程数超过GPU能同时执行的数目后,反而会因为调度开销而变慢。一般建议线程数是SM数的几十倍,而不是几万倍。
- 误区二:忽略内存访问模式
GPU对连续内存访问非常友好。如果你让线程随机访问内存,性能会暴跌。我有个项目,就因为没注意内存对齐,性能差了10倍。
- 误区三:在GPU里写复杂分支
前面说了,分支发散是性能杀手。尽量把复杂逻辑放到CPU上,GPU只做简单的、可预测的计算。
核心总结:从CPU到GPU,不是简单的“多开几个线程”。而是要从“任务并行”思维切换到“数据并行”思维。你要学会把问题分解成大量独立的小任务,让成千上万个线程同时处理。
7. 动手试试:第一个GPU程序
最后,我建议你写一个最简单的GPU程序:给数组每个元素加1。
__global__ void add_one(int* data, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
data[idx] += 1;
}
}
int main() {
int n = 1024;
int* d_data;
cudaMalloc(&d_data, n * sizeof(int));
// ... 拷贝数据、启动核函数、拷贝结果
add_one<<<1, 256>>>(d_data, n);
// ...
}
这个程序虽然简单,但它包含了GPU编程的核心思想:每个线程处理一个元素,所有线程执行相同的操作。
嗯,等你跑通了这个程序,你就迈出了从CPU思维到GPU思维的第一步。下一章,我们会深入讨论线程层次结构——Grid、Block、Thread之间的关系。
记住:别用CPU的思维写GPU代码。把循环忘掉,拥抱线程。