第二个CUDA程序：Hello World内核与主机设备代码分离

说实话，学CUDA编程最让人兴奋的，就是亲手写出第一个能在GPU上跑的内核。我记得当年第一次看到屏幕上打出"Hello World"时，心里那个激动——虽然只是个简单的打印，但背后是CPU和GPU协同工作的完整流程。

这一章，我们就来写这个经典的入门程序。不过我不会只给你看代码，我会把背后的编译流程、API选择这些坑，都一一说清楚。

1. 一个完整的Hello World内核

先看代码，再慢慢拆解。

// hello.cu
#include <stdio.h>

// 这是GPU上运行的kernel函数
__global__ void helloFromGPU()
{
    printf("Hello World from GPU! blockIdx.x = %d, threadIdx.x = %d\n",
           blockIdx.x, threadIdx.x);
}

int main()
{
    // 主机端代码
    printf("Hello World from CPU!\n");

    // 启动kernel：1个block，每个block里5个线程
    helloFromGPU<<<1, 5>>>();

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

    return 0;
}

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

__global__：这个修饰符告诉编译器，这个函数要在GPU上跑，并且可以从CPU调用。我刚开始写的时候，经常忘记加这个，结果编译报错半天找不到原因。
<<<1, 5>>>：这是CUDA特有的语法，叫"执行配置"。第一个参数是block数量，第二个参数是每个block里的线程数。说白了，就是告诉GPU：你给我开1个工作组，每个组里5个工人干活。
cudaDeviceSynchronize()：这个函数让CPU等着，直到GPU把活干完。我曾经踩过一个坑——不加这行，CPU直接跑完了，程序退出，GPU上的printf还没输出呢。

2. 主机代码与设备代码分离

你可能会问：为什么要把CPU代码和GPU代码写在一起？

其实，这是CUDA编程模型的核心思想——异构计算。CPU负责控制逻辑、I/O这些串行任务，GPU负责大规模并行计算。代码混在一起，但编译时会分开处理。

我个人习惯把kernel函数单独放在一个文件里，比如kernel.cu，主机代码放在main.cpp里。这样项目结构清晰，也方便复用。

关键概念：主机（Host）指CPU及其内存，设备（Device）指GPU及其显存。两者通过PCIe总线通信，数据传输是开销最大的操作之一。

3. nvcc编译流程——到底发生了什么？

你写了一个.cu文件，然后执行nvcc hello.cu -o hello，就生成了可执行文件。但背后发生了什么？

我画了一张图，帮你理清整个流程。

说白了，nvcc就是个"翻译官"。它把.cu文件拆成两部分：GPU代码交给它自己的编译器，CPU代码丢给系统自带的C++编译器（比如g++或cl.exe）。最后再把两边的结果拼起来。

小技巧：你可以用 nvcc --keep 选项保留中间文件。这样能看到生成的PTX代码和C++中间文件，对理解编译过程很有帮助。我调试复杂kernel时经常用这招。

4. 运行时API vs 驱动API

CUDA提供了两套API来跟GPU打交道。你猜怎么着？我们刚才用的cudaDeviceSynchronize()就是运行时API的函数。

特性	运行时API (Runtime API)	驱动API (Driver API)
易用性	高，函数封装好，上手快	低，需要手动管理更多细节
初始化	自动初始化，无需显式调用	需要 `cuInit()` 手动初始化
错误处理	返回 `cudaError_t`	返回 `CUresult`
模块管理	编译时链接，简单	支持运行时加载 `.cubin` 或 `.ptx`
适用场景	大多数应用开发	需要精细控制、动态加载的场景

我个人建议：初学者先用运行时API。它封装得好，代码简洁，不容易出错。等你把CUDA玩熟了，再去看驱动API——它给你更多控制权，但代价是代码量翻倍。

避坑指南：我曾经在一个项目里混用了两套API，结果出了个诡异的bug——运行时API自动初始化了上下文，驱动API又手动初始化了一次，导致上下文冲突。排查了两天才找到原因。所以，一个程序里只用一套API，别混着来。

5. 动手试试：编译和运行

把上面的代码保存为hello.cu，然后执行：

nvcc hello.cu -o hello
./hello

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

Hello World from CPU!
Hello World from GPU! blockIdx.x = 0, threadIdx.x = 0
Hello World from GPU! blockIdx.x = 0, threadIdx.x = 1
Hello World from GPU! blockIdx.x = 0, threadIdx.x = 2
Hello World from GPU! blockIdx.x = 0, threadIdx.x = 3
Hello World from GPU! blockIdx.x = 0, threadIdx.x = 4

看到了吗？CPU打印了一次，GPU打印了5次——因为我们在配置里写了<<<1, 5>>>，启动了5个线程。每个线程都执行了helloFromGPU函数，只是threadIdx.x不同。

你想想看，如果我把配置改成<<<2, 3>>>，会输出几行？答案是6行——2个block，每个block里3个线程。这就是CUDA最基本的并行模型：网格(Grid) → 块(Block) → 线程(Thread)。

嗯，这一章的内容就到这里。代码虽短，但背后的编译流程、API选择这些知识，是后续所有CUDA编程的基础。把这些搞明白了，后面的路就好走了。

想要更多资料，联系可可

📚 课程目录

第二个CUDA程序：Hello World内核与主机设备代码分离

1. 一个完整的Hello World内核

2. 主机代码与设备代码分离

3. nvcc编译流程——到底发生了什么？

4. 运行时API vs 驱动API

5. 动手试试：编译和运行