第一个CUDA程序：Hello World
好，咱们直接动手写代码。说实话，学CUDA最怕的就是光看不练。我见过太多人把理论背得滚瓜烂熟，一开编辑器就懵了。今天咱们就写一个最简单的CUDA程序，让你亲眼看到GPU在干活。
先理解两个世界：主机与设备
在CUDA的世界里，有两个角色：主机（Host）和设备（Device）。主机就是你的CPU和内存，设备就是GPU和它的显存。说白了，CPU负责调度，GPU负责猛算。
我刚开始学的时候，总搞混一个概念：CPU和GPU的内存是不互通的。你不能直接在GPU上访问CPU的变量，反过来也不行。数据得先拷贝过去，算完了再拷回来。这个坑我踩过好几次，后面会细说。
| 角色 | 硬件 | 内存 | 职责 |
|---|---|---|---|
| 主机（Host） | CPU | 系统内存（DDR） | 控制逻辑、串行任务 |
| 设备（Device） | GPU | 显存（VRAM） | 大规模并行计算 |
写一个Hello World
先别急着写并行代码。咱们从最基础的开始——让GPU说句话。
// hello.cu
#include <stdio.h>
// 这是运行在GPU上的函数
__global__ void helloFromGPU()
{
printf("Hello World from GPU!\n");
}
int main()
{
// 调用GPU函数，<<<1, 1>>>表示用1个线程块，每个块1个线程
helloFromGPU<<<1, 1>>>();
// 同步等待GPU执行完毕
cudaDeviceSynchronize();
printf("Hello World from CPU!\n");
return 0;
}注意看那个 __global__ 关键字。它告诉编译器：这个函数要在GPU上跑，并且可以从CPU调用。我个人习惯把这种函数叫做「内核函数」或「Kernel」。
还有那个奇怪的 <<<1, 1>>> 语法。这是CUDA特有的调用方式，尖括号里第一个数字是线程块数量，第二个是每个块里的线程数。这里咱们只启动1个线程，所以只打印一次。
关键点：所有GPU上的printf输出，必须等 cudaDeviceSynchronize() 之后才能看到。不然CPU跑完了，GPU还没打印完，你就看不到输出。
编译与运行
CUDA程序用 nvcc 编译器。它不是gcc，但底层会调用gcc来编译主机代码。
# 编译
nvcc -o hello hello.cu
# 运行
./hello如果你看到这样的输出：
Hello World from GPU!
Hello World from CPU!恭喜你，第一个CUDA程序跑通了！
嗯，这里要注意：GPU打印的顺序不一定在CPU之前。因为GPU是异步执行的，CPU的printf可能先执行。但加了 cudaDeviceSynchronize() 之后，GPU的打印一定会先完成。
常见错误排查
写CUDA程序，90%的时间都在调试。我把自己踩过的坑整理一下，你遇到了直接对照着查。
错误1：nvcc找不到
如果你输入 nvcc 提示命令未找到，说明CUDA Toolkit没装或者环境变量没配。检查一下 /usr/local/cuda/bin 有没有加到PATH里。
我的习惯：在 ~/.bashrc 里加一行 export PATH=/usr/local/cuda/bin:$PATH，然后 source ~/.bashrc。
错误2：编译报错 "error: identifier 'printf' is undefined"
这是因为在GPU代码里用printf，需要包含 <stdio.h>。但注意，不是所有GPU都支持printf。计算能力2.0以上的设备才支持。你可以用 cudaGetDeviceProperties 查一下。
错误3：运行时卡死或没输出
最常见的原因：忘了加 cudaDeviceSynchronize()。我曾经有一次调试了半小时，发现GPU其实跑完了，但CPU直接退出了，输出被冲掉了。
警告：千万别在生产代码里滥用 cudaDeviceSynchronize()。它会阻塞CPU，让并行变串行。调试时用用可以，正式跑的时候要删掉。
错误4：段错误（Segmentation Fault）
这通常是因为你访问了非法内存。比如在GPU代码里用了CPU的指针。记住：GPU函数里只能访问 cudaMalloc 分配的内存，或者用 cudaMemcpy 拷贝过来的数据。
让程序跑得更「并行」一点
刚才那个程序只用了1个线程，太浪费GPU了。咱们改成启动多个线程试试。
__global__ void helloFromGPU()
{
// threadIdx.x 是当前线程在块内的编号
printf("Hello from thread %d\n", threadIdx.x);
}
int main()
{
// 启动10个线程
helloFromGPU<<<1, 10>>>();
cudaDeviceSynchronize();
return 0;
}运行后你会看到10行输出，但顺序是乱的。为什么？因为GPU的线程是并行执行的，谁先跑完谁先打印。你想想看，如果10个人同时喊话，你听到的顺序肯定不是固定的。
核心概念：GPU线程的执行顺序是不确定的。你不能依赖线程间的执行顺序来做逻辑判断。如果需要同步，得用 __syncthreads()，这个后面会讲。
调试小技巧
我刚开始写CUDA时，最头疼的就是调试。CPU程序可以用gdb单步，GPU程序怎么单步？后来发现几个实用方法：
- 用printf大法：在GPU代码里加printf，打印关键变量的值。虽然慢，但最直接。
- 检查返回值：每个CUDA API调用都会返回一个错误码。养成习惯，每次调用后检查
cudaGetLastError()。 - 用cuda-memcheck：运行程序时加上
cuda-memcheck ./hello，它会帮你检测内存越界、未初始化变量等问题。
避坑指南：我曾经在项目里遇到一个诡异的bug——程序有时跑得对，有时跑得错。查了两天，最后发现是某个变量没初始化。GPU的全局变量默认是0，但局部变量不是。所以记得初始化所有变量。
总结一下
今天咱们干了三件事：
- 写了一个在GPU上打印Hello World的程序
- 学会了用nvcc编译和运行
- 知道了常见的坑怎么避开
说实话，这个程序虽然简单，但它包含了CUDA编程的所有核心要素：主机与设备的区分、内核函数的定义、线程的启动方式、以及同步机制。后面的所有内容，都是在这个基础上加东西。
下一章咱们会深入线程模型，看看GPU到底是怎么同时跑成千上万个线程的。到时候你会发现，刚才的10个线程只是开胃菜。