3、第一个CUDA程序：Hello World核函数、编译与运行、错误检查

好，咱们正式开始写代码了。

说实话，我当年学CUDA的时候，最怕的就是环境配好了，第一个程序跑不通。那种挫败感，你懂的。所以这一章，我带你写一个最简的CUDA程序——说白了就是GPU版的"Hello World"。别小看它，这里面藏着核函数、编译流程、错误处理三大核心知识点。

3.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;
}

嗯，这里有几个关键点，我一个个说。

3.1.1 __global__ 修饰符

看到 __global__ 了吗？这是CUDA的"身份证"。它告诉编译器：这个函数要在GPU上跑，而且可以从CPU端调用。

我个人习惯把 __global__ 函数叫做"入口函数"。它就像一扇门，CPU通过这扇门把任务交给GPU。

3.1.2 三括号语法 <<<1, 10>>>

这个三括号，是CUDA最标志性的语法。它用来配置核函数的执行参数。

• 第一个参数：网格中线程块的数量（这里是1个块）
• 第二个参数：每个线程块中的线程数（这里是10个线程）

所以 <<<1, 10>>> 的意思是：启动1个线程块，每个块里有10个线程。总共10个线程并行执行 helloFromGPU。

我在项目中遇到过一个问题：新手容易把这两个参数搞反。记住，第一个是块数，第二个是每块的线程数。你想想看，如果写成了 <<<10, 1>>>，那就是10个块，每个块1个线程，结果虽然也是10个线程，但线程的组织结构完全不同——这会影响后续的线程索引计算。

3.1.3 cudaDeviceSynchronize()

为什么需要这个？

因为CPU和GPU是异步工作的。CPU调用 helloFromGPU<<<1, 10>>>() 后，不会等GPU执行完，而是直接继续往下走。如果没有 cudaDeviceSynchronize()，程序可能直接结束，你根本看不到GPU打印的内容。

3.2 编译与运行

写好了代码，怎么跑？

3.2.1 使用nvcc编译器

CUDA的编译器叫 nvcc。它会把核函数编译成GPU能执行的PTX代码，同时把主机代码交给宿主编译器（比如gcc）处理。

编译命令很简单：

nvcc -o hello hello.cu

这里 .cu 是CUDA源文件的扩展名。-o hello 指定输出文件名为 hello。

运行：

./hello

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

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

为什么会打印10次？因为启动了10个线程，每个线程都执行了一次printf。

3.2.2 常见编译错误

我曾经因为忘记加 cudaDeviceSynchronize()，调试了整整一个下午。输出时有时无，我还以为是显卡坏了……嗯，从那以后我再也不敢省略同步操作了。

3.3 错误检查：写CUDA程序的必修课

CUDA API函数大部分返回 cudaError_t 类型。如果你不检查返回值，出错了都不知道。

3.3.1 手动检查模式

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

这样写太啰嗦了。我一般用宏封装：

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

然后就可以这样用：

CHECK(cudaDeviceSynchronize());

是不是清爽多了？

3.3.2 核函数内部的错误

注意：核函数本身不返回值。你无法用 CHECK() 直接检查核函数调用。正确的做法是：

helloFromGPU<<<1, 10>>>();
CHECK(cudaGetLastError());  // 检查核函数启动是否成功
CHECK(cudaDeviceSynchronize());  // 检查执行过程中是否有错误

cudaGetLastError() 会返回最近一次异步调用的错误。如果核函数启动参数非法（比如块数太多），这里就能抓到。

3.4 本章知识体系

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

3.5 完整示例：带错误检查的Hello World

最后，给你一个可以直接复制使用的完整版本：

#include <stdio.h>
#include <cuda_runtime.h>

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

__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(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());

    printf("All done!\n");
    return 0;
}

编译运行：

nvcc -o hello hello.cu
./hello

你会看到每个GPU线程打印自己的 threadIdx.x，从0到9。这就是线程索引——下一章我们会深入讲这个。

好了，第一个CUDA程序就到这里。你学会了写核函数、用nvcc编译、做错误检查。这三板斧，够你应付大部分入门场景了。

想要更多资料，联系可可