CUDA编程模型：主机与设备、核函数、线程层次结构

大家好，我是你们的老朋友。今天咱们聊聊CUDA编程模型的核心三要素。说实话，我刚接触CUDA那会儿，也被这些概念绕得头晕。但后来我发现，只要抓住几个关键点，这东西其实挺直观的。

你想想看，CUDA编程模型说白了就是回答三个问题：谁干活？怎么分配活？活怎么干？ 嗯，咱们一个一个来拆解。

主机与设备：CPU和GPU的分工

在CUDA的世界里，主机（Host）指的是CPU及其内存系统，设备（Device）指的是GPU及其显存。这两者各司其职，配合好了效率翻倍。

我个人习惯把CPU想象成项目经理，GPU就是一群埋头苦干的工人。项目经理负责统筹、调度、处理复杂逻辑；工人负责重复性的体力活——大规模并行计算。

我在项目中遇到过不少新手，恨不得把所有计算都扔到GPU上。其实没必要。比如一个简单的if-else分支判断，CPU处理比GPU快得多。GPU的优势在于「同一件事做一万遍」，而不是「一万件不同的事各做一遍」。

主机和设备之间通过PCIe总线通信。这里有个坑——数据传输很贵。我曾经优化过一个图像处理程序，发现80%的时间都花在CPU和GPU之间来回拷贝数据上。所以我的建议是：尽量减少数据传输次数，能一次传完就别分十次。

核函数（Kernel）：GPU上执行的函数

核函数，就是跑在GPU上的函数。用__global__修饰符声明，调用时用尖括号语法<<<grid, block>>>指定执行配置。

来看个最简单的例子：

// 核函数声明
__global__ void vecAdd(float* A, float* B, float* C, int n) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

// 主机端调用
int main() {
    // ... 数据准备和拷贝 ...
    vecAdd<<<numBlocks, blockSize>>>(d_A, d_B, d_C, n);
    // ... 结果回传 ...
}

注意几个关键点：

• 核函数必须返回void，不能有返回值
• 核函数内部不能调用主机端函数（比如printf在计算能力2.0以上才支持）
• 核函数是异步执行的——调用后CPU会继续执行下一行代码，不会等GPU算完

说到异步，我踩过一个大坑。有一次我核函数调用完直接就去读结果，读出来的全是垃圾数据。后来才想起来忘了加cudaDeviceSynchronize()。嗯，这里要注意：核函数调用后，如果需要立即使用结果，必须显式同步。

线程层次结构：Grid、Block、Thread

这是CUDA最精妙的设计之一。说白了就是三级组织架构：

• Thread（线程）：最小的执行单元，每个线程执行核函数的一份拷贝
• Block（线程块）：一组线程的集合，块内线程可以共享内存、同步协作
• Grid（网格）：一组线程块的集合，一个核函数对应一个Grid

为什么搞这么复杂？直接弄一堆线程不行吗？

原因在于GPU的硬件架构。GPU由多个SM（流多处理器）组成，每个SM可以调度多个Block。Block内的线程运行在同一个SM上，可以高效协作。而不同Block之间是独立的，不能直接通信。

线程的索引通过内置变量获取：

• threadIdx.x：线程在块内的索引
• blockIdx.x：块在网格内的索引
• blockDim.x：每个块的线程数
• gridDim.x：网格内的块数

全局线程ID的计算公式：int tid = threadIdx.x + blockIdx.x * blockDim.x;

这个公式我建议你刻在脑子里。我见过太多bug都是因为索引算错了——要么越界，要么多个线程操作同一个数据。

知识体系总览

下面这张图把今天讲的内容串起来了：

实际编程中的注意事项

我曾经接手过一个项目，对方把Block大小设成了32，结果GPU利用率只有可怜的15%。我改成256后，性能直接翻了4倍。你想想看，同样的代码，只是改了个参数，效果天差地别。

最后说一句：理解线程层次结构是CUDA优化的基石。Grid、Block、Thread这三层关系搞明白了，后面学共享内存、同步、流处理都会轻松很多。别急，慢慢来，多写几个demo跑一跑，感觉自然就来了。

想要更多资料，联系可可