GPU架构与核函数：CUDA核心概念

说实话，刚接触CUDA那会儿，我也被一堆新名词搞得头晕。Grid、Block、Thread，还有各种内存——全局内存、共享内存、寄存器……这些东西到底怎么配合？今天我就带你捋一捋。

GPU不是CPU的简单替代品。它是另一种思考方式。我习惯把GPU想象成一个超级工厂，里面有成千上万个工人同时干活。嗯，这个比喻虽然老套，但确实管用。

核心思想：GPU通过大量轻量级线程并行执行，隐藏内存访问延迟，实现高吞吐量计算。

线程层次结构：Grid / Block / Thread

CUDA把线程组织成三层结构。说白了就是：

Thread（线程）——最基础的执行单元。每个线程执行同一个核函数，但处理不同的数据。
Block（线程块）——一组线程的集合。同一个Block里的线程可以协作，共享数据。
Grid（网格）——所有Block的集合。一个Grid对应一次核函数调用。

我在项目中遇到过一个问题：Block大小设成多少最合适？当时我试了各种组合，最后发现128或256个线程/Block通常表现不错。为什么？因为GPU的调度单元（Warp）一次处理32个线程，128正好是4个Warp，调度起来很顺畅。

个人经验：Block大小建议设为32的倍数。128、256、512都是常见选择。太小了隐藏不了延迟，太大了寄存器压力大。

来看看代码怎么写：

// 核函数定义
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    // 计算全局线程ID
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// 调用核函数
int main() {
    int N = 1 << 20;  // 约100万个元素
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
}

你想想看，这里blockIdx.x是Block的索引，threadIdx.x是线程在Block里的索引，blockDim.x是Block的大小。三者一组合，就得到了全局唯一的线程ID。

避坑指南：我曾经犯过一个低级错误——忘记检查线程ID是否越界。当数据量不是Block大小的整数倍时，多余的线程会访问非法内存。所以一定要加 if (idx < N) 这个判断。

内存层次结构

GPU的内存不是铁板一块。它有多个层级，速度和容量各不相同。我画了张图帮你理解：

这张图你看懂了吗？越靠近核心（SM），速度越快，但容量越小。我刚开始写CUDA程序时，总喜欢把所有数据放全局内存，结果性能惨不忍睹。后来才明白——数据放对地方，比算法优化更重要。

各内存类型详解

内存类型	作用域	生命周期	速度	典型用途
全局内存	所有线程 + Host	程序运行期间	慢（~400周期）	存储大量数据
共享内存	同一Block内线程	Block生命周期	快（~20周期）	线程间数据共享、归约
寄存器	单个线程	线程生命周期	极快（1周期）	局部变量、中间结果
常量内存	所有线程（只读）	程序运行期间	快（有缓存）	常量参数、查找表
局部内存	单个线程	线程生命周期	慢（同全局内存）	寄存器溢出时的后备

关键点：共享内存是程序员可以控制的缓存。用好它，性能能翻好几倍。我做过一个矩阵乘法的优化，用共享内存后速度提升了8倍。

共享内存的使用示例

来看一个实际例子。矩阵乘法中，每个输出元素需要读取一整行和一整列的数据。如果不加优化，全局内存访问次数多到吓人。

__global__ void matMulShared(float *A, float *B, float *C, int N) {
    // 声明共享内存
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    // 分块计算
    for (int t = 0; t < N / TILE_SIZE; t++) {
        // 协作加载数据到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        
        __syncthreads();  // 同步：确保所有线程加载完成
        
        // 从共享内存计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        
        __syncthreads();  // 同步：确保所有线程计算完成
    }
    
    C[row * N + col] = sum;
}

这里有个细节：__syncthreads()是个同步屏障。同一Block里的所有线程必须都执行到这里，才能继续往下走。我曾经漏掉一个__syncthreads()，结果数据还没加载完就开始计算，算出来的结果全是错的。嗯，调试这种bug特别痛苦。

我的习惯：每次使用共享内存后，都问自己三个问题——数据加载完了吗？同步了吗？有没有bank conflict？

Warp与线程调度

GPU执行的最小单位不是线程，而是Warp。一个Warp包含32个线程，它们一起执行同一条指令。这就是所谓的SIMT（单指令多线程）模型。

你想想看，如果32个线程里有16个走if分支，16个走else分支，会发生什么？两个分支都会被执行，但只有一半的线程在工作。这就是Warp divergence（分支发散），性能杀手。

避坑指南：我曾经写过一个核函数，里面有个if-else判断线程ID的奇偶性。结果性能直接腰斩。后来改成用位运算统一处理，速度就回来了。记住：让同一个Warp里的线程走相同的执行路径。

内存合并访问

全局内存访问有个重要原则：连续线程访问连续地址。当Warp里的32个线程访问的内存地址是连续的，GPU会把它们合并成一次或几次大的内存事务。否则，就是多次小事务，效率极低。

举个例子：

// 好的访问模式：合并访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];  // 线程0访问data[0]，线程1访问data[1]...

// 坏的访问模式：非合并访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx * 64];  // 线程0访问data[0]，线程1访问data[64]...

我在项目中遇到过这种情况：一个看起来没问题的算法，跑起来就是慢。后来用NVIDIA的Nsight工具一分析，发现内存访问效率只有20%。改成合并访问后，效率直接拉到95%以上。

总结一下：GPU编程的核心就是管理好两样东西——线程的协作方式和数据的存放位置。线程层次结构决定了协作范围，内存层次结构决定了数据访问速度。把这两样搞明白了，你的核函数就成功了一大半。

好了，这一章的内容就到这里。记住我上面说的那些坑，写代码时多留个心眼。下一章我们会聊聊更具体的优化技巧，到时候见。

想要更多资料，联系可可

📚 课程目录

GPU架构与核函数：CUDA核心概念

线程层次结构：Grid / Block / Thread

内存层次结构

各内存类型详解

共享内存的使用示例

Warp与线程调度

内存合并访问