4、第一个CUDA程序：从Hello World开始

说实话，学任何编程语言，第一个程序都是Hello World。CUDA也不例外。但CUDA的Hello World有点不一样——它不是在屏幕上打印一句话那么简单，而是让你第一次感受到GPU并行计算的魅力。

我个人习惯，学新东西先跑通一个最小示例。这样心里有底，后面再深入也不慌。好，我们直接上手。

4.1 编写你的第一个Kernel

先看代码，再解释。下面这个程序，会在GPU上打印"Hello World from GPU!"，每个线程打印一次。

#include <stdio.h>

// 定义kernel函数
__global__ void helloFromGPU()
{
    printf("Hello World from GPU! threadId: %d\n", threadIdx.x);
}

int main()
{
    // 启动kernel，1个block，10个线程
    helloFromGPU<<<1, 10>>>();

    // 同步，确保GPU执行完毕
    cudaDeviceSynchronize();

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

注意看，__global__这个修饰符。它告诉编译器：这个函数要在GPU上执行，但由CPU调用。嗯，这里有个小坑——我刚开始写的时候，总忘记加这个修饰符，结果编译报错，找了半天原因。

尖括号里的<<<1, 10>>>是CUDA特有的语法。第一个数字是block数量，第二个是每个block里的线程数。说白了，就是告诉GPU：你给我开10个线程，跑这个函数。

4.2 nvcc编译流程

CUDA程序不能用普通的gcc编译。你得用NVIDIA提供的编译器——nvcc。

编译上面那个程序，命令很简单：

nvcc hello.cu -o hello

然后运行：

./hello

你会看到类似这样的输出：

Hello World from GPU! threadId: 0
Hello World from GPU! threadId: 1
...
Hello World from GPU! threadId: 9
Hello World from CPU!

nvcc到底干了什么？我简单说一下它的工作流程：

1. 分离代码：把.cu文件中的GPU代码（带__global__、__device__的部分）和CPU代码分开
2. 编译GPU代码：用NVIDIA的编译器把GPU代码编译成PTX（一种中间表示）或者直接编译成cubin（二进制）
3. 编译CPU代码：调用宿主机的C++编译器（比如g++）编译CPU部分
4. 链接：把两部分链接成最终的可执行文件

4.3 运行时API vs 驱动API

CUDA提供了两套API让你跟GPU打交道。很多初学者搞不清区别，我当年也迷糊过。

对比项	运行时API (Runtime API)	驱动API (Driver API)
易用性	简单，上手快	复杂，需要更多代码
初始化	自动初始化	需要手动cuInit()
模块管理	编译时链接	支持运行时加载.cubin
函数前缀	cuda开头	cu开头
使用场景	大部分应用	需要精细控制时

说白了，运行时API是对驱动API的封装。它帮你处理了很多琐事，比如设备初始化、上下文管理。我个人建议：初学者先用运行时API，等你对CUDA熟悉了，再考虑驱动API。

什么时候需要用驱动API？我在项目中遇到过一种情况：需要动态加载不同的GPU内核，根据输入数据决定用哪个。这时候驱动API的cuModuleLoad就派上用场了。运行时API做不到这点。

4.4 错误检查机制

写CUDA程序，最怕什么？GPU出错了，但CPU程序还在跑，结果数据全错，你还不知道。我刚开始学的时候，就吃过这个亏。

CUDA的API函数大部分会返回一个错误码。但很多人（包括曾经的我）会忽略它。正确的做法是：每次调用CUDA API后，都检查返回值。

我习惯写一个宏来做这件事：

#define CHECK(call)                                   \
do {                                                  \
    const cudaError_t err = call;                     \
    if (err != cudaSuccess) {                         \
        printf("CUDA Error: %s at %s:%d\n",           \
               cudaGetErrorString(err),               \
               __FILE__, __LINE__);                   \
        exit(EXIT_FAILURE);                           \
    }                                                 \
} while(0)

// 使用示例
CHECK(cudaMalloc(&d_data, size));
CHECK(cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice));
CHECK(cudaDeviceSynchronize());

我曾经在一个项目里，忘了检查cudaMemcpy的返回值。结果数据拷贝失败了，但程序继续跑，最后算出来的结果全是错的。排查了整整两天，才发现是这个问题。从那以后，我每个CUDA API调用都加上了错误检查。

4.5 本章知识体系

下面这张图，帮你理清本章的核心逻辑：

嗯，这张图把本章的四个核心知识点串起来了。你写第一个CUDA程序时，就按这个思路来：先写Kernel，再用nvcc编译，搞清楚用哪套API，最后别忘了加错误检查。

想要更多资料，联系可可