第2章 CUDA编程模型

2. CUDA编程模型：线程层次结构与内存层次结构

好，咱们今天聊聊CUDA里最核心的两个概念——线程层次结构和内存层次结构。说实话，这两个东西搞明白了，你写CUDA程序就成功了一半。我当年刚接触CUDA时，就是被这两层结构绕晕过，后来踩了不少坑才慢慢摸清楚门道。

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

先说说线程是怎么组织的。CUDA把线程分成了三层：Grid（网格）→ Block（线程块）→ Thread（线程）。你可以想象成一个公司：Grid是总公司，Block是部门，Thread就是每个员工。

为什么这么设计？说白了，是为了适配GPU的硬件架构。GPU里有多个SM（流多处理器），每个SM可以同时跑多个Block。你想想看，如果只有一层线程结构，那调度起来得多麻烦？

核心关系：

一个Grid包含多个Block
一个Block包含多个Thread
同一个Block里的Thread可以共享数据（通过共享内存）
不同Block的Thread之间不能直接通信

来看个代码示例，这是我常用的一个核函数模板：

// 典型的CUDA核函数
__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    // 计算全局线程ID
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查
    if (tid < N) {
        C[tid] = A[tid] + B[tid];
    }
}

这里blockIdx.x是Block在Grid中的索引，threadIdx.x是Thread在Block中的索引，blockDim.x是Block里线程的数量。这个公式我建议你记牢了，几乎每个CUDA程序都会用到。

2.2 内存层次结构

说完线程，咱们聊聊内存。CUDA的内存层次比线程还复杂，我画了个图帮你理解：

嗯，这里要注意：全局内存是所有线程都能读写的，但速度最慢。共享内存是同一个Block里的线程共享的，速度比全局内存快很多，但容量有限。我习惯把共享内存当成「程序员手动管理的缓存」来用。

避坑指南：

我曾经在项目里犯过一个低级错误——把频繁访问的数据放在全局内存里，结果性能惨不忍睹。后来改成用共享内存做数据复用，速度直接提升了5倍。记住：能用共享内存就别用全局内存，能用寄存器就别用共享内存。

2.3 内存访问模式：合并访问

说到内存，不得不提一个关键概念——合并访问（Coalesced Access）。什么意思呢？就是相邻的线程最好访问相邻的内存地址。GPU在读取全局内存时，会一次性读取128字节（一个cache line）。如果线程0读地址0，线程1读地址4，线程2读地址8...那一次读取就能满足所有线程，效率极高。

反过来，如果线程0读地址0，线程1读地址1000，线程2读地址2000...那每次读取只能服务一个线程，效率就低得可怜。

// 合并访问（高效）
__global__ void coalescedAccess(float *data, int N) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) {
        data[tid] = data[tid] * 2.0f;  // 相邻线程访问相邻地址
    }
}

// 非合并访问（低效）
__global__ void stridedAccess(float *data, int N, int stride) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) {
        data[tid * stride] = data[tid * stride] * 2.0f;  // 线程间有间隔
    }
}

警告：非合并访问是CUDA性能杀手之一。我见过不少新手写的程序，明明算法没问题，就是慢，一查发现是内存访问模式不对。你想想看，GPU有几千个线程同时跑，如果每个线程都要等内存，那性能能好才怪。

2.4 实际项目中的选择策略

那在实际项目中，怎么选择Block大小和内存类型呢？我一般遵循这几个原则：

场景	推荐Block大小	内存策略	原因
数据并行（如向量加）	256-512线程	全局内存直接访问	计算简单，无数据复用
矩阵乘法	16x16 或 32x32	共享内存做分块	数据复用率高，减少全局内存访问
归约操作	256线程	共享内存+同步	需要线程间协作
图像处理	16x16	纹理内存	空间局部性好，硬件加速

我个人习惯是：Block大小选256或512，因为这样能保证每个SM上有足够的线程来隐藏内存延迟。太小了（比如32）会导致SM利用率低，太大了（比如1024）又可能受限于寄存器数量。

核心总结：

线程层次：Grid → Block → Thread，决定了并行粒度
内存层次：全局内存 → 共享内存 → 寄存器，决定了访问速度
合并访问：相邻线程访问相邻地址，这是性能优化的基本功
Block大小：256-512是黄金区间，具体看算法需求

好了，这一章的内容就这些。记住，CUDA编程模型的核心就是「怎么组织线程」和「怎么管理内存」。这两件事搞明白了，后面学优化技巧就轻松多了。