第一章：踏上GPU编程之路——从Hello World到内存管理

各位同学，欢迎来到《GPU算子开发能力提升》的第一课。今天我们要聊的内容，可以说是GPU编程的“敲门砖”。你可能会想：不就是个Hello World吗？有什么好讲的？

嗯，我当年也是这么想的。直到我在一个项目里，因为cudaMalloc分配的内存大小不对，排查了整整两天……从那以后，我再也不敢小看这些基础操作了。

1.1 第一个CUDA程序：让GPU说Hello

先来看一个最经典的例子。我们让GPU打印一句话：

#include <stdio.h>

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

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

这段代码里，__global__ 是关键。它告诉编译器：这个函数要在GPU上运行。而 <<<1, 10>>> 是CUDA的启动配置，表示启动1个线程块，每个块里10个线程。

核心要点： __global__ 修饰的函数必须返回void，且只能由CPU调用。

我个人习惯在写第一个CUDA程序时，先确认GPU设备是否可用。你可以用 cudaGetDeviceCount() 来检查。有一次我在新服务器上跑代码，发现怎么都打印不出来，结果发现是驱动没装好……

1.2 向量加法：GPU并行计算的入门案例

Hello World只是开胃菜。真正体现GPU价值的是并行计算。我们来看向量加法：

__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

这里有个公式你要记住：全局线程ID = 块ID × 块大小 + 线程ID。说白了，就是每个线程负责计算一个元素。

你想想看，如果N=100万，用CPU得循环100万次。但GPU可以同时启动成千上万个线程，一次搞定。这就是并行计算的魅力。

小技巧： 启动配置时，建议让线程块大小是32的倍数（比如128、256）。因为GPU的warp大小就是32，这样能充分利用硬件资源。

1.3 内存管理：cudaMalloc / cudaMemcpy / cudaFree

GPU不能直接访问CPU的内存。你得先把数据拷贝过去。这就涉及三个核心函数：

函数	作用	注意事项
cudaMalloc	在GPU上分配内存	返回的是设备指针，不能直接在CPU上解引用
cudaMemcpy	在CPU和GPU之间拷贝数据	需要指定方向：cudaMemcpyHostToDevice 或 cudaMemcpyDeviceToHost
cudaFree	释放GPU内存	忘记释放会导致内存泄漏，严重时程序崩溃

来看一个完整的向量加法示例：

int main() {
    int N = 1000;
    size_t size = N * sizeof(float);

    // 1. 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_A[i] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 3. 拷贝数据到GPU
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 4. 启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 5. 拷贝结果回CPU
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 6. 释放内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

避坑指南： 我曾经在项目里忘记调用cudaFree，结果程序跑了几个小时突然崩溃。排查后发现是显存泄漏。记住：每次cudaMalloc都要对应一个cudaFree，就像malloc和free一样。

1.4 错误检查与调试：别让bug过夜

CUDA编程最让人头疼的是什么？是错误信息不够直观。很多时候，核函数出错了，程序直接崩溃，连个提示都没有。

所以我强烈建议你封装一个错误检查宏：

#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(&d_A, size));
CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
CHECK(cudaDeviceSynchronize());  // 检查核函数执行错误

为什么要在核函数后面加 cudaDeviceSynchronize()？因为核函数是异步执行的。CPU不会等GPU算完就继续往下走。如果你不同步，后面的错误检查可能捕获不到核函数里的问题。

调试技巧： 用 cuda-memcheck 工具可以检测越界访问、未初始化内存等问题。我每次写完核函数，都会跑一遍这个工具，能省去很多排查时间。

1.5 知识体系总览

下面这张图总结了本章的核心知识点，你可以把它当作学习路线图：

这张图把四个核心模块串起来了。从左到右，从最简单的Hello World开始，到向量加法理解并行，再到内存管理掌握数据流动，最后用错误检查保证代码质量。每一步都是下一章的基础。

好了，第一章的内容就到这里。记住：写CUDA代码，先保证能跑通，再追求性能。别一上来就想着优化，先把基础打牢。

想要更多资料，联系可可

📚 课程目录

第一章：踏上GPU编程之路——从Hello World到内存管理

1.1 第一个CUDA程序：让GPU说Hello

1.2 向量加法：GPU并行计算的入门案例

1.3 内存管理：cudaMalloc / cudaMemcpy / cudaFree

1.4 错误检查与调试：别让bug过夜

1.5 知识体系总览