4. CUDA编程基础回顾：线程模型与内存层次

好，咱们今天来聊聊CUDA编程最核心的两个东西——线程模型和内存层次。这两个概念，说白了就是你在GPU上写代码的「世界观」。我刚开始学CUDA的时候，总觉得这玩意儿不就是多线程嘛，有啥好学的？结果一上手就踩坑，跑出来的结果全是错的。嗯，后来才明白，不理解线程怎么组织和内存怎么分布，你写的Kernel大概率是废的。

4.1 CUDA线程模型：Grid、Block、Thread

先说说线程模型。GPU上跑的程序叫Kernel，一个Kernel启动时，会创建成千上万个线程。这些线程不是乱糟糟的，而是有组织有纪律的。

我习惯这么理解：Thread是最小的工作单元，多个Thread组成一个Block，多个Block组成一个Grid。Grid对应一个Kernel启动，Block对应一个SM（流多处理器）上的执行单元，Thread就是实际干活的人。

关键点：一个Block内的Thread可以协作（通过共享内存和同步），不同Block的Thread之间不能直接通信。这是CUDA编程的黄金法则。

每个Thread和Block都有唯一的ID。你写代码的时候，经常要用到这些ID来索引数据。比如：

// 获取当前线程在全局中的索引
int tid = blockIdx.x * blockDim.x + threadIdx.x;

这段代码我几乎每个Kernel都会写。blockIdx是Block的索引，threadIdx是Thread在Block内的索引，blockDim是Block的维度大小。三者一乘一加，就得到了全局唯一的线程ID。

你想想看，如果你要处理一个长度为N的数组，每个线程处理一个元素，那这个tid就是数组的下标。简单吧？但坑也在这里——边界检查一定要做。我曾经有一次忘了检查，结果线程数比数组长度多，读到了脏数据，排查了整整一下午。

我的习惯：启动Kernel时，Grid和Block的维度通常设为一维，除非处理图像这种二维数据。一维最简单，不容易出错。

4.2 内存层次结构：Global、Shared、Local

讲完线程，咱们聊聊内存。GPU的内存层次，说白了就是「谁离计算单元近，谁就快，但容量小」。我画了一张图，帮你理清这个结构：

CUDA内存层次结构 Global Memory（全局内存） 容量大（GB级） | 延迟高（~400 cycles） | 所有线程可访问 Shared Memory（共享内存） 容量小（几十KB） | 延迟低（~30 cycles） | 同一Block内线程共享 Local Memory（局部内存） 每个线程私有 | 实际存在Global中 寄存器 最快 | 容量极少 慢 快

这张图从上到下，速度越来越快，容量越来越小。咱们一个一个说。

4.1.1 Global Memory（全局内存）

这是GPU上最大的内存，所有线程都能读写。但它的延迟很高，大概400个时钟周期。你想想看，如果每个线程都频繁访问Global Memory，那性能肯定上不去。

避坑指南：我曾经写过一个矩阵乘法的Kernel，直接拿Global Memory做累加，结果比CPU还慢。后来改成Shared Memory做分块，速度直接翻了10倍。所以，能用Shared就别用Global，这是优化第一原则。

4.1.2 Shared Memory（共享内存）

Shared Memory是每个Block私有的，Block内的所有线程都能访问。它的延迟只有30个时钟周期左右，比Global快了一个数量级。但容量很小，一般只有几十KB。

Shared Memory的典型用法是做数据复用。比如矩阵乘法，你把一个子矩阵加载到Shared Memory里，然后Block内的所有线程反复用它，这样就避免了反复访问Global Memory。

注意：Shared Memory使用不当会导致Bank Conflict（存储体冲突），这会降低访问效率。我刚开始写的时候完全没注意这个，后来用NVIDIA Nsight分析才发现性能瓶颈在这里。解决办法是让线程访问的地址尽量分散到不同的Bank。

4.1.3 Local Memory（局部内存）

Local Memory听起来像是「局部」的，但其实它物理上存在Global Memory里。每个线程有自己的Local Memory，用于存放寄存器放不下的变量（比如大数组）。

嗯，这里要注意：Local Memory的延迟和Global一样高。所以如果你的Kernel里声明了一个大数组，编译器可能会把它放到Local Memory里，性能就会下降。我建议尽量用寄存器变量，或者显式使用Shared Memory。

4.3 简单的Kernel编写与Launch

理论讲完了，咱们写个简单的Kernel练练手。就拿向量加法来说吧，这是CUDA的「Hello World」。

// 向量加法的Kernel
__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];
    }
}

// 主函数中启动Kernel
int main() {
    int N = 1 << 20;  // 约100万个元素
    size_t size = N * sizeof(float);
    
    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);
    
    // 初始化数据...
    
    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // 拷贝数据到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    
    // 配置Grid和Block
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    // 启动Kernel
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // 拷贝结果回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    
    // 清理...
    return 0;
}

这段代码有几个关键点：

  • __global__ 关键字告诉编译器这是一个Kernel函数，在设备上执行，从主机端调用。
  • threadsPerBlock 我习惯设成256，这是经验值。128也行，但256通常能更好地利用SM的资源。
  • blocksPerGrid 的计算公式 (N + threadsPerBlock - 1) / threadsPerBlock 是向上取整，确保覆盖所有元素。
  • 边界检查 if (i < N) 必不可少。因为线程数可能比元素多，多出来的线程什么都不做。

我的建议：写Kernel时，先把功能跑通，再考虑优化。我见过太多人一上来就搞Shared Memory、搞循环展开，结果代码跑出来全是错的。先保证正确性，再用工具分析瓶颈，最后优化。

启动Kernel时用的三括号语法 <<<grid, block>>> 是CUDA特有的。第一个参数是Grid的维度（Block的数量），第二个参数是Block的维度（Thread的数量）。这两个参数都可以是一维、二维或三维的。处理图像时，我经常用二维Grid和二维Block，这样索引起来更直观。

好了，这一章的内容就这些。线程模型和内存层次是CUDA编程的基石，理解透了，后面的优化技巧才能用得上。下一章咱们聊聊更高级的优化手段，比如如何用Shared Memory做矩阵乘法加速。

本章核心回顾：

  • Grid → Block → Thread 三级层次，Block内线程可协作
  • Global Memory容量大但慢，Shared Memory快但小
  • Local Memory是寄存器的后备，尽量少用大数组
  • 写Kernel先保证正确，再优化性能

想要更多资料，联系可可

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