第一个CUDA程序：Hello World从CPU到GPU

说实话，我教GPU编程这么多年，发现很多新手卡在第一关——搞不清CPU和GPU到底怎么分工。你想想看，写个普通的C程序，main函数里跑得欢，突然要扔一部分代码到显卡上跑，这思维转换确实需要点时间。

我个人习惯把GPU编程比作「外包」。CPU是项目经理，GPU是外包团队。项目经理把任务拆好、数据打包，喊一声「干活了」，GPU团队闷头算完，再把结果交回来。嗯，这个比喻虽然糙，但道理不糙。

CUDA程序的基本骨架

先看一个最简的CUDA程序。别怕，代码量很少，我带你一行行拆。

#include <stdio.h>

// 设备端代码 —— 跑在GPU上
__global__ void helloFromGPU()
{
    printf("Hello World from GPU!\n");
}

// 主机端代码 —— 跑在CPU上
int main()
{
    // 调用GPU内核函数
    helloFromGPU<<<1, 1>>>();

    // 同步等待GPU执行完毕
    cudaDeviceSynchronize();

    printf("Hello World from CPU!\n");
    return 0;
}

这段代码你编译运行后，会先看到GPU打印的消息，再看到CPU打印的消息。为什么顺序是这样？因为cudaDeviceSynchronize()强制CPU等着GPU干完活。

主机端 vs 设备端：代码分离的艺术

上面代码里，main()是主机端代码，跑在CPU上。helloFromGPU前面有个__global__，这是设备端代码，跑在GPU上。

我在项目中遇到过不少新手，把printf直接写在__global__函数里，然后抱怨「怎么没输出？」——GPU的printf是异步的，你不加cudaDeviceSynchronize()，程序可能已经退出了，输出缓冲区还没刷新呢。

代码位置	运行硬件	关键字/语法
主机端 (Host)	CPU	普通C/C++函数
设备端 (Device)	GPU	__global__ 修饰的函数

__global__ 关键字：GPU的入口函数

__global__是CUDA里最重要的关键字之一。它告诉编译器：这个函数要被编译成GPU指令，并且可以从CPU端调用。

有几个硬性规则，我建议你记牢：

• __global__函数必须返回void。为什么？因为GPU内核启动是异步的，你没法直接拿到返回值。想传数据回来？用指针参数。
• 调用时必须用<<<...>>>语法，不能像普通函数那样直接写func()。
• 函数体内可以调用__device__函数（跑在GPU上的辅助函数），但不能调用__host__函数（CPU函数）。

<<< >>> 执行配置语法：告诉GPU怎么干活

这个三括号语法是CUDA独有的。它的格式是：

kernel_name<<<gridDim, blockDim, sharedMem, stream>>>(args...);

前两个参数是必填的：

• gridDim：线程网格的维度。可以是一个整数，也可以是dim3类型（三维）。
• blockDim：每个线程块的维度。同样可以是整数或dim3。

后两个参数可选：

• sharedMem：每个线程块分配的共享内存字节数，默认0。
• stream：指定在哪个CUDA流上执行，默认0（默认流）。

咱们刚才的例子<<<1, 1>>>，意思就是：启动1个线程块，每个线程块里有1个线程。总共就1个线程在干活，说白了就是单线程跑GPU，纯粹为了演示。

知识体系一览

下面这张图帮你理清本章的核心脉络。我画图时习惯把「主机-设备分离」放在最上面，因为这是理解一切的基础。

编译与运行

保存代码为hello.cu，然后用nvcc编译：

nvcc hello.cu -o hello
./hello

如果你用的是HIP（AMD GPU），命令类似：

hipcc hello.cu -o hello
./hello

避坑指南

• 忘记同步：GPU内核是异步的，不加cudaDeviceSynchronize()，CPU可能先跑完退出，GPU输出还没打印。
• 文件名后缀：CUDA源文件必须用.cu后缀，否则nvcc不认。
• GPU不支持printf：计算能力1.x的老卡不支持设备端printf。不过现在市面上基本见不到了，放心。

好了，第一个CUDA程序就讲到这里。你把这个例子跑通，CUDA编程的大门就算正式打开了。下一章咱们聊聊线程层次结构——说白了就是怎么让成千上万个线程有序地干活。

---

公众号:蓝海资料掘金营，微信deep3321