第一章：第一个CUDA程序——Hello World核函数

大家好，我是你们的老朋友。今天咱们聊聊CUDA编程的起点——怎么写第一个核函数。

说实话，我第一次接触CUDA时，心里也犯嘀咕：这玩意儿跟普通C++到底差在哪？后来踩了几个坑才明白，核函数就是跑在GPU上的函数，但写法上有些门道。

1.1 核函数长什么样？

先看一个最简单的例子。你想想看，CPU上打印"Hello World"多简单，一行printf搞定。但在GPU上，你得告诉它：每个线程都要打印一次。

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

__global__ void helloFromGPU() {
    printf("Hello World from GPU!\n");
}

int main() {
    helloFromGPU<<<1, 10>>>();
    cudaDeviceSynchronize();
    return 0;
}

注意那个 __global__ 修饰符。这是CUDA的关键字，告诉编译器：这个函数要跑在GPU上，由CPU调用。我在项目中见过有人忘了加这个，结果函数在CPU上跑，性能完全不对。

<<<1, 10>>> 是执行配置。第一个参数是线程块数，第二个是每个块里的线程数。这里我启动了1个块，每个块10个线程，总共10个线程并行执行。

1.2 nvcc编译流程——别被吓到

CUDA代码不能直接用gcc编译。你得用nvcc。说白了，nvcc是个编译器驱动，它会把代码拆成两部分：GPU代码交给NVCC编译器，CPU代码交给gcc/clang。

编译命令很简单：

nvcc -o hello hello.cu

如果你有多个文件，可以这样：

nvcc -c hello.cu -o hello.o
nvcc -c main.cpp -o main.o
nvcc hello.o main.o -o program

我曾经犯过一个低级错误：在Makefile里忘了加 -arch=sm_75，结果程序在V100上跑得慢得要命。后来才发现，nvcc默认生成的是最低兼容的代码，性能损失很大。

nvcc -arch=sm_80 -o hello hello.cu  # 针对Ampere架构

1.3 运行时API——你真正打交道的东西

CUDA运行时API是一组C函数，让你在CPU端管理GPU。常用的有：

嗯，这里要注意：cudaMalloc 返回的是设备指针，你不能在CPU上直接解引用。我见过新手直接 printf("%d", *devPtr)，结果段错误。

1.4 错误检查——别让bug悄悄溜走

CUDA API调用失败时，会返回错误码。但很多人不检查。我刚开始也这样，直到有一次 cudaMalloc 失败，程序继续跑，结果数据全乱了。

我的习惯是写一个宏：

#define CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", \
                    __FILE__, __LINE__, \
                    cudaGetErrorString(err)); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

// 使用示例
CHECK(cudaMalloc(&devPtr, size));
CHECK(cudaMemcpy(devPtr, hostPtr, size, cudaMemcpyHostToDevice));

1.5 完整示例——把知识串起来

咱们写一个完整的Hello World，带错误检查：

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

#define CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "Error at %s:%d: %s\n", \
                    __FILE__, __LINE__, \
                    cudaGetErrorString(err)); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

__global__ void helloFromGPU() {
    printf("Hello from thread %d in block %d\n", 
           threadIdx.x, blockIdx.x);
}

int main() {
    // 获取GPU信息
    cudaDeviceProp prop;
    CHECK(cudaGetDeviceProperties(&prop, 0));
    printf("GPU: %s\n", prop.name);
    
    // 启动核函数
    helloFromGPU<<<2, 5>>>();
    CHECK(cudaDeviceSynchronize());
    
    printf("All threads completed!\n");
    return 0;
}

编译运行：

nvcc -arch=sm_80 -o hello hello.cu
./hello

输出类似：

GPU: NVIDIA A100
Hello from thread 0 in block 0
Hello from thread 1 in block 0
...
Hello from thread 4 in block 1

1.6 知识体系图

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

这张图展示了从编写核函数到最终执行的完整流程。你想想看，每一步都环环相扣。少了错误检查，bug可能藏得很深；忘了指定架构，性能就上不去。

好了，第一章就到这里。记住：写CUDA程序，核心就是三件事——写核函数、管理内存、检查错误。把这三点吃透，后面的路就好走了。