4. 第一个CUDA程序：Hello World from GPU

说实话，每次学一门新技术，我第一件事就是跑个Hello World。这就像跟新硬件打个招呼——嘿，能听见我说话吗？CUDA也不例外。今天我们就从最简单的GPU程序开始，看看它到底怎么工作的。

4.1 核函数（Kernel）是什么？

核函数，说白了就是在GPU上执行的函数。它跟CPU上的普通函数最大的区别是——它会被成千上万个线程同时执行。

我刚开始接触CUDA时，总觉得核函数很神秘。后来发现，它其实就是加了个__global__修饰符的C函数。嗯，就这么简单。

核函数的三个特征：

用__global__声明，告诉编译器这是给GPU跑的
返回类型必须是void，不能有返回值
调用时要用特殊的<<>>语法

4.2 第一个GPU程序：Hello World

来，直接上代码。这是我当年写的第一个CUDA程序，现在看还挺亲切的。

#include <stdio.h>

// 核函数定义：每个线程打印自己的ID
__global__ void helloFromGPU() {
    printf("Hello World from GPU! Thread ID: %d\n", threadIdx.x);
}

int main() {
    // 调用核函数：启动4个线程
    helloFromGPU<<<1, 4>>>();
    
    // 等待GPU执行完毕
    cudaDeviceSynchronize();
    
    printf("Hello World from CPU!\n");
    return 0;
}

编译命令：

nvcc hello.cu -o hello
./hello

输出结果：

Hello World from GPU! Thread ID: 0
Hello World from GPU! Thread ID: 1
Hello World from GPU! Thread ID: 2
Hello World from GPU! Thread ID: 3
Hello World from CPU!

看到了吗？GPU上的4个线程几乎同时执行，每个线程都打印了自己的ID。这就是并行计算最直观的体现。

4.3 <<>> 语法详解

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

其实它的含义很直接：

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

gridDim：线程网格的维度，即有多少个线程块
blockDim：线程块的维度，即每个块有多少个线程

总线程数 = gridDim × blockDim

我的经验：刚开始写CUDA时，我总搞混grid和block。后来我这样记——grid是「大部队」，block是「小分队」。每个小分队里有若干士兵（线程）。

4.4 主机与设备交互

CPU叫主机（Host），GPU叫设备（Device）。它们之间怎么通信？我画了张图，一看就明白。

主机和设备交互的核心就是数据搬运。我总结了一个三步走套路：

准备数据：在主机内存里把数据准备好
搬运数据：用cudaMemcpy把数据拷到显存
启动核函数：GPU干活，结果再拷回来

注意：我曾经犯过一个低级错误——核函数启动后没加cudaDeviceSynchronize()，结果CPU先跑完了，程序直接退出，GPU上的活儿全白干了。记住，核函数是异步的，CPU不会等它。

4.5 完整的交互示例

来个实际点的例子，在GPU上做向量加法：

#include <stdio.h>

// 核函数：向量加法
__global__ void vecAdd(float *A, float *B, float *C, int n) {
    int i = threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    int n = 5;
    float h_A[] = {1, 2, 3, 4, 5};
    float h_B[] = {10, 20, 30, 40, 50};
    float h_C[5];
    
    // 1. 在显存上分配空间
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, n * sizeof(float));
    cudaMalloc(&d_B, n * sizeof(float));
    cudaMalloc(&d_C, n * sizeof(float));
    
    // 2. 把数据从主机拷到设备
    cudaMemcpy(d_A, h_A, n * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, n * sizeof(float), cudaMemcpyHostToDevice);
    
    // 3. 启动核函数
    vecAdd<<<1, n>>>(d_A, d_B, d_C, n);
    
    // 4. 把结果拷回主机
    cudaMemcpy(h_C, d_C, n * sizeof(float), cudaMemcpyDeviceToHost);
    
    // 5. 打印结果
    for (int i = 0; i < n; i++) {
        printf("%.0f + %.0f = %.0f\n", h_A[i], h_B[i], h_C[i]);
    }
    
    // 6. 释放显存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    
    return 0;
}

输出：

1 + 10 = 11
2 + 20 = 22
3 + 30 = 33
4 + 40 = 44
5 + 50 = 55

避坑指南：我曾经忘了cudaFree，结果跑了几百次后显存爆了。GPU不像CPU有虚拟内存，显存用完就崩。养成好习惯——malloc和free成对出现。

4.6 本章小结

这一章我们干了三件事：

写了第一个CUDA程序，让GPU说了句Hello World
搞懂了<<>>语法——grid是块数，block是每块的线程数
学会了主机和设备怎么配合——数据搬运三步走

说实话，CUDA入门最大的门槛就是理解「并行」这个概念。CPU是一个人在干活，GPU是一万个人同时干活。你想想看，分配任务的方式能一样吗？

下一章我们会深入聊聊线程的组织方式。嗯，到时候你就知道为什么要有grid和block这两层结构了。

📚 课程目录

4. 第一个CUDA程序：Hello World from GPU

4.1 核函数（Kernel）是什么？

4.2 第一个GPU程序：Hello World

4.3 <<>> 语法详解

4.4 主机与设备交互

4.5 完整的交互示例

4.6 本章小结