4. 第一个CUDA程序：编写Hello World核函数、编译与运行、理解<<<>>>语法、错误处理

好，咱们直接动手写代码。

学任何编程语言，第一个程序都是 Hello World。CUDA 也不例外。不过，CUDA 的 Hello World 不是打印一句话那么简单——它要让你看到 GPU 在干活。

4.1 编写第一个核函数

先看代码。我习惯把核函数写在前面，这样逻辑清晰。

#include <stdio.h>

// 核函数定义
__global__ void helloFromGPU(void)
{
    printf("Hello World from GPU!\n");
}

int main(void)
{
    // 主机端打印
    printf("Hello World from CPU!\n");

    // 启动核函数
    helloFromGPU<<<1, 10>>>();

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

    return 0;
}

嗯，这里要注意几个关键点。

第一，__global__ 这个修饰符。 它告诉编译器：这个函数是在 GPU 上跑的，但要从 CPU 端调用。说白了，这就是 CPU 和 GPU 之间的桥梁。

第二，printf 能在核函数里用？ 是的，从 CUDA 2.0 开始支持。我在项目中经常用它做快速调试，比用 cuda-gdb 快多了。不过正式代码里别这么干，性能会崩。

4.2 理解 <<<>>> 语法

这个三括号语法，是 CUDA 最标志性的东西。我第一次看到它时，心想：这什么鬼？

其实它很简单——就是告诉 GPU 要开多少个线程。

格式是：

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

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

拿上面的例子来说：

helloFromGPU<<<1, 10>>>();

这表示：启动 1 个线程块，每个块里有 10 个线程。总共 10 个线程在跑。

你想想看，如果改成 <<<2, 5>>>，结果一样是 10 个线程。但组织方式不同——2 个块，每块 5 个线程。这在后面讲线程协作时会很重要。

4.3 编译与运行

编译 CUDA 程序，用 nvcc 编译器。它是 NVIDIA 的 CUDA 编译器，会把核函数编译成 GPU 能执行的 PTX 或二进制代码。

nvcc -o hello hello.cu

然后运行：

./hello

输出应该是：

Hello World from CPU!
Hello World from GPU!
Hello World from GPU!
...（共10次）

4.4 错误处理——别偷懒

CUDA 编程里，错误处理是重中之重。我刚开始学的时候，经常核函数调用失败，程序直接崩了，连个错误信息都没有。

为什么会这样？因为核函数启动是异步的。CPU 不会等 GPU 执行完就继续往下跑。如果核函数出错了，CPU 根本不知道。

所以，我们需要显式检查错误。

最常用的方法是检查 cudaGetLastError() 的返回值：

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
    return -1;
}

我建议把它封装成一个宏，方便复用：

#define CHECK(call) \
{ \
    const cudaError_t err = call; \
    if (err != cudaSuccess) { \
        printf("Error: %s:%d, ", __FILE__, __LINE__); \
        printf("code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
        exit(1); \
    } \
}

然后这样用：

CHECK(cudaDeviceSynchronize());

4.5 完整示例与执行流程

把上面的代码整合一下，加上错误处理：

#include <stdio.h>

#define CHECK(call) \
{ \
    const cudaError_t err = call; \
    if (err != cudaSuccess) { \
        printf("Error: %s:%d, ", __FILE__, __LINE__); \
        printf("code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
        exit(1); \
    } \
}

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

int main(void)
{
    printf("Hello World from CPU!\n");

    helloFromGPU<<<1, 10>>>();

    CHECK(cudaDeviceSynchronize());

    return 0;
}

这里我加上了 threadIdx.x，它是 CUDA 内置变量，表示当前线程在块内的索引。从 0 到 9，正好对应 10 个线程。

整个执行流程是这样的：

1. CPU 执行 main 函数，打印 "Hello World from CPU!"
2. 遇到 <<<1, 10>>>，CPU 把核函数发射到 GPU
3. GPU 创建 10 个线程，每个线程独立执行 helloFromGPU
4. 每个线程打印自己的 threadIdx.x 值
5. cudaDeviceSynchronize() 等待所有 GPU 线程完成
6. 程序退出

下面这张图展示了这个流程：

这张图把整个流程串起来了。你看，CPU 负责启动和同步，GPU 负责并行执行。分工明确。

好了，第一个 CUDA 程序就到这里。代码不多，但背后的概念很重要——线程、块、网格、同步、错误处理。这些是后续所有内容的基础。

---

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