2. CUDA编程基础：主机与设备、核函数编写、线程层次结构

好，咱们正式开始动手写CUDA代码了。这一章是地基中的地基，你把它啃透了，后面那些花里胡哨的优化技巧才能玩得转。我当年刚接触CUDA时，就是没搞懂主机和设备到底怎么分工，结果调试一个简单的向量加法调了整整两天——后来发现是数据根本没拷到显卡上。嗯，这种坑咱们今天必须绕过去。

2.1 主机与设备：CPU和GPU的分工

说白了，CUDA程序里有两个角色：主机（Host）和设备（Device）。主机就是CPU及其内存，设备就是GPU及其显存。CPU擅长逻辑控制和串行计算，GPU擅长大规模并行计算。所以我们的策略是：CPU负责调度、控制流程、管理数据；GPU负责闷头算数。

核心原则：数据在哪里，计算就在哪里。你要在GPU上算，就得先把数据搬到GPU的显存里。

我习惯把主机和设备想象成两个独立的仓库。CPU仓库（主机内存）里的货，GPU仓库（显存）拿不到。你得叫一辆卡车（cudaMemcpy）把货拉过去。算完了，再拉回来。就这么简单。

来看一下典型的数据流：

  1. 在主机内存中分配空间并初始化数据
  2. 在设备显存中分配空间
  3. 将数据从主机拷贝到设备
  4. 在GPU上执行核函数
  5. 将结果从设备拷贝回主机
  6. 释放主机和设备上的内存

注意：忘记释放显存是新手最常见的错误。显存泄漏比内存泄漏更隐蔽，因为你的电脑可能不会立刻崩溃，但跑几次大模型训练后，显卡就罢工了。我曾经在生产环境里排查过这种问题，最后发现是某个中间结果没释放，导致OOM。所以，养成好习惯：cudaMalloc 和 cudaFree 成对出现。

2.2 核函数（Kernel）编写

核函数就是在GPU上执行的函数。它的写法跟普通C函数很像，但有几个关键区别。你想想看，如果每个线程都执行同一个函数，那它们怎么知道自己该干哪一部分活？答案是通过内置变量。

先看一个最简单的核函数——向量加法：

// 核函数定义：前面加 __global__ 修饰符
__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];
    }
}

这里有几个要点：

  • __global__ 告诉编译器，这个函数是在设备上执行，但从主机端调用
  • 核函数返回类型必须是 void
  • 核函数不能使用C++的STL容器、不能有静态变量、不能有可变参数
  • 核函数内部可以调用设备函数（用 __device__ 修饰）

我的习惯：写核函数时，我总会在开头先算好线程索引，然后立刻做边界检查。为什么？因为线程总数不一定能被块大小整除，最后一批线程可能会越界。不做检查，轻则读到垃圾数据，重则程序崩溃。这个习惯帮我避免过无数次bug。

调用核函数的方式也很特别：

// 主机端调用核函数
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

尖括号里的两个参数，第一个是网格大小（多少个线程块），第二个是块大小（每个块多少个线程）。这两个参数决定了你总共启动多少个线程。

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

这是CUDA最核心的概念之一。GPU的线程不是一维的，而是三维的层次结构。我刚开始学的时候也觉得三维有点多余，但后来做图像处理时才发现，用二维线程索引访问像素简直不要太爽。

来看这张图，它展示了线程的组织方式：

Grid（网格） Block (0,0) Block (1,0) Block (2,0) Block (0,1) Block (1,1) Block (2,1) T0 T1 T2 T3 T4 T5 T6 T7 T8 T9 T10 T11 T12 T13 T14 T15 T16 T17 Block Thread 一个Grid包含多个Block，每个Block包含多个Thread

每个线程都有三个内置变量来确定自己的位置：

内置变量 类型 说明
threadIdx dim3 线程在块内的索引（x, y, z）
blockIdx dim3 块在网格内的索引（x, y, z）
blockDim dim3 每个块的维度大小（有多少个线程）
gridDim dim3 网格的维度大小（有多少个块）

计算全局线程索引的公式是：

// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 二维情况（常用于图像处理）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;  // 展平为一维

关于块大小的选择：我个人建议块大小取32的倍数，最好是128或256。为什么是32？因为GPU的warp大小就是32，一个warp是GPU调度的最小单位。块大小取32的倍数，能保证warp里没有空闲线程。我踩过的坑是：有一次为了省事设了100，结果每个warp里只有4个线程干活，28个在摸鱼，性能直接腰斩。

2.4 完整的向量加法示例

把上面所有知识点串起来，写一个完整的程序：

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

// 核函数：向量加法
__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];
    }
}

int main() {
    int N = 1 << 20;  // 1048576 个元素
    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] = 1.0f;
        h_B[i] = 2.0f;
    }
    
    // 2. 设备端分配显存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // 3. 拷贝数据到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    
    // 4. 配置线程并启动核函数
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;  // 向上取整
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    
    // 5. 拷贝结果回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }
    
    // 6. 释放内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);
    
    return 0;
}

避坑指南：我曾经在项目里忘记加 cudaDeviceSynchronize()，结果核函数还没执行完，主机就开始读结果，读到的全是0。核函数是异步的，主机不会等它执行完。如果你需要在核函数之后立即读取结果，记得加同步。不过在这个例子里，cudaMemcpy 本身会隐式同步，所以没问题。

2.5 小结

这一章我们搞清楚了三个核心概念：主机和设备的分工、核函数的写法、线程的三层组织结构。说白了，CUDA编程就是：CPU负责指挥和搬运，GPU负责干活；每个线程通过内置变量找到自己该处理的数据；线程按Grid-Block-Thread三层组织，方便我们处理各种维度的数据。

嗯，这些概念刚接触时可能有点绕，但写两三个小例子就熟了。我建议你现在就动手把上面的向量加法跑一遍，然后试着改成矩阵加法（二维线程）。代码跑通了，这些概念就真正是你的了。

想要更多资料，联系可可

← 上一章 📖 返回目录 下一章 →