3. CUDA编程入门：从零开始理解GPU编程

好，咱们直接进入正题。CUDA编程，说白了就是让CPU和GPU一起干活。CPU擅长复杂逻辑，GPU擅长并行计算。怎么让它们配合？这就是CUDA要解决的问题。

我个人习惯把CUDA编程模型理解成「主从模式」。CPU是主控，GPU是计算工人。CPU发号施令，GPU埋头苦干。这个思路贯穿整个CUDA编程，你记住它，后面就好理解了。

3.1 CUDA编程模型：CPU与GPU的分工

CUDA的编程模型其实不复杂。我画个简单的图给你看：

• Host（主机）：CPU + 系统内存。负责控制流程、数据准备、启动核函数。
• Device（设备）：GPU + 显存。负责大规模并行计算。

程序跑起来是这样的流程：

1. CPU准备数据，放到系统内存里。
2. CPU调用 cudaMalloc 在显存上分配空间。
3. CPU调用 cudaMemcpy 把数据从系统内存拷贝到显存。
4. CPU启动核函数（Kernel），GPU开始并行计算。
5. 计算完成后，CPU再把结果从显存拷回来。

嗯，这里要注意：CPU和GPU是异步工作的。CPU启动核函数后，不会傻等GPU算完，它会继续执行后面的代码。如果你需要等GPU算完，得手动加同步。

3.2 核函数（Kernel）编写：GPU上跑的函数长什么样

核函数，就是跑在GPU上的函数。写法上有个标志：__global__ 修饰符。

举个例子，一个简单的向量加法：

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

// 调用核函数
int main() {
    // ... 数据准备和内存分配 ...
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    // ... 结果拷贝 ...
}

看到没？核函数有几个特点：

• 返回类型必须是 void。GPU函数不能有返回值。
• 用 <<<grid, block>>> 语法启动。这俩参数决定了有多少线程干活。
• 函数体里用 threadIdx、blockIdx 这些内置变量来定位当前线程。

我在项目中遇到过一个问题：新手容易把核函数写成普通函数，忘了加 __global__。结果编译过了，但跑起来全是0。排查了半天才发现是核函数根本没在GPU上执行。嗯，这种坑踩过一次就记住了。

3.3 线程层次结构：Grid、Block、Thread

这是CUDA里最核心的概念。我刚开始学的时候也觉得绕，后来想明白了一个比喻：

• Thread（线程）：一个工人，干一份活。
• Block（线程块）：一个小组，组内工人可以互相通信（共享内存）。
• Grid（线程网格）：整个工厂，包含所有小组。

为什么要有这个层次？说白了是为了适配GPU的硬件架构。GPU里有多个流多处理器（SM），每个SM可以跑多个Block。Block里的线程跑在同一个SM上，所以可以共享数据。

看个具体的配置例子：

// 假设有 10000 个元素要处理
int N = 10000;
int blockSize = 256;          // 每个Block 256个线程
int gridSize = (N + blockSize - 1) / blockSize;  // 向上取整

vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

这里 blockSize 我习惯设成256。为什么？因为这是经验值。太小了SM利用率低，太大了寄存器不够用。256是个比较平衡的选择。

线程索引的计算公式要记牢：

int tid = threadIdx.x + blockIdx.x * blockDim.x;

这个公式我每次写核函数都要默念一遍。说白了就是：当前线程在全局中的位置 = 线程在块内的位置 + 块在网格中的位置 × 块的大小。

3.4 内存管理：cudaMalloc 和 cudaMemcpy

GPU不能直接访问CPU的内存。你得先把数据搬到显存里。这就涉及两个核心API：

看个完整的内存管理示例：

float *h_A, *h_B, *h_C;  // CPU端指针
float *d_A, *d_B, *d_C;  // GPU端指针
int N = 10000;
size_t size = N * sizeof(float);

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

// 2. 分配GPU显存
cudaMalloc(&d_A, size);
cudaMalloc(&d_B, size);
cudaMalloc(&d_C, size);

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

// 4. 启动核函数
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

// 5. 结果从GPU拷回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);

这里有个细节：cudaMemcpy 的最后一个参数指定方向。有四种：

• cudaMemcpyHostToDevice：CPU → GPU
• cudaMemcpyDeviceToHost：GPU → CPU
• cudaMemcpyDeviceToDevice：GPU → GPU（显存内部拷贝）
• cudaMemcpyHostToHost：CPU → CPU（基本不用）

#define CHECK(call) { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        printf("Error: %s\n", cudaGetErrorString(err)); \
        exit(1); \
    } \
}

嗯，说到内存管理，我曾经犯过一个低级错误：核函数跑完了，忘了 cudaMemcpy 把结果拷回来，直接打印CPU端的数组。结果全是0。排查了半小时才发现数据还在显存里没动。你想想看，这种错误多冤啊。

3.5 小结

这一章的内容，说白了就是CUDA编程的「三板斧」：

1. 写核函数：用 __global__ 修饰，用 <<<grid, block>>> 启动。
2. 定线程结构：Grid、Block、Thread 三层，用公式算好索引。
3. 管好内存：cudaMalloc 分配，cudaMemcpy 搬运，cudaFree 释放。

这三步走顺了，CUDA编程就算入门了。下一章咱们聊聊「如何写出高效的核函数」，包括内存访问模式、共享内存的使用这些实战技巧。到时候我会分享一些我在实际项目中踩过的坑，保证让你少走弯路。