CUDA编程基础回顾：GPU内存层次结构、Kernel函数编写、线程块与网格配置、共享内存使用、同步机制

各位同学，欢迎来到《TensorCore实战指南》的第一章。

说实话，每次带新人入门CUDA，我都会先问一个问题：你清楚数据在GPU里是怎么流动的吗？ 很多人上来就写kernel，结果性能惨不忍睹。嗯，这其实不怪你——GPU的内存模型确实比CPU复杂得多。今天我们就从最底层开始，把这块硬骨头啃下来。

本章核心脉络：理解GPU内存层次 → 学会写kernel → 合理配置线程 → 用好共享内存 → 掌握同步。

1. GPU内存层次结构：从慢到快，从大到小

GPU的内存，说白了就是一个金字塔结构。越往上走，速度越快，容量越小。我刚开始做CUDA优化时，总想着把所有数据都塞进共享内存，结果发现根本装不下——这就是没理解层次结构的代价。

内存类型	位置	访问延迟	作用域	生命周期
全局内存	片外DRAM	~400-800周期	所有线程	程序运行期间
共享内存	片上SRAM	~5-30周期	同一线程块	线程块生命周期
寄存器	片上	~1周期	单个线程	线程生命周期
常量内存	片外(带缓存)	~5-10周期(缓存命中)	所有线程	程序运行期间
纹理内存	片外(带缓存)	~5-10周期(缓存命中)	所有线程	程序运行期间

个人经验：我一般把全局内存比作「仓库」，共享内存比作「工作台」。数据从仓库搬到工作台上处理，效率最高。千万别在仓库里直接干活——那太慢了。

2. Kernel函数编写：从Hello World到实战

写一个CUDA kernel，其实就是在告诉GPU：「嘿，每个线程都去执行这段代码」。但这里有个坑——你写的代码是给一个线程看的，不是给所有线程看的。我见过不少新手把串行逻辑直接搬进kernel，结果跑出来全是错的。

来看一个最简单的向量加法：

// 一个线程处理一个元素
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

注意那个 if (idx < N) 边界检查。为什么需要它？因为你的线程总数可能比数据量多。我早期有一次忘了加这个判断，结果数组越界，查了整整一个下午的bug。

避坑指南：我曾经在写kernel时，把 threadIdx.x 和 blockIdx.x 搞混了。记住：threadIdx 是线程在块内的索引，blockIdx 是块在网格内的索引。两者相乘才能得到全局索引。

3. 线程块与网格配置：找到那个甜蜜点

配置线程块和网格，说白了就是决定「派多少人、分几个组去干活」。这里有个关键参数——warp。GPU以32个线程为一组（一个warp）来调度。所以你的block大小最好是32的倍数。

我个人习惯这样配置：

block大小：128或256。太小了（比如32）会导致SM利用率低；太大了（比如1024）又可能受限于寄存器数量。
grid大小：根据数据量计算，保证覆盖所有元素。
维度选择：处理1D数据用1D grid+1D block；处理2D图像用2D grid+2D block。

// 典型的配置方式
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

为什么block大小是256比较好？ 因为每个SM能同时驻留的线程数有限（比如1024个线程）。用256的block，一个SM可以塞4个block，正好填满。用128的话，得8个block才能填满，调度开销更大。

4. 共享内存使用：片上缓存的艺术

共享内存是CUDA优化的核心武器。它位于芯片上，速度比全局内存快一个数量级。但它的容量很小——通常每个SM只有几十KB到一百多KB。

共享内存的典型用法：

把全局内存中的数据加载到共享内存（协作加载）
在共享内存中做计算（快速访问）
把结果写回全局内存

// 矩阵转置示例：使用共享内存
__global__ void transpose(float *in, float *out, int width) {
    __shared__ float tile[TILE_DIM][TILE_DIM];
    
    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;
    
    // 协作加载到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 转置后写回
    int x_out = blockIdx.y * TILE_DIM + threadIdx.x;
    int y_out = blockIdx.x * TILE_DIM + threadIdx.y;
    out[y_out * width + x_out] = tile[threadIdx.x][threadIdx.y];
}

避坑指南：共享内存有个叫「bank conflict」的东西。简单说，如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。我曾在做矩阵乘法时，因为没注意bank conflict，性能直接腰斩。解决办法？调整数据布局，或者用padding。

5. 同步机制：别让线程跑乱了

同步是CUDA里最容易出bug的地方。你想想看，几千个线程同时跑，如果没有同步机制，数据读写顺序就全乱了。

CUDA提供了几种同步方式：

__syncthreads()：块内屏障。所有线程必须到达这个点才能继续。常用于共享内存读写之间。
原子操作：比如 atomicAdd，用于全局求和、计数等场景。
内存栅栏：__threadfence()，保证内存操作的可见性。

我曾经踩过的坑：在if-else分支里用了__syncthreads()。结果某些线程进了if，某些进了else，永远等不到对方——死锁了。记住：__syncthreads()必须被块内所有线程执行到，不能放在条件分支里。

原子操作虽然好用，但性能开销大。我一般只在必须保证顺序的场景下才用。比如直方图统计、全局归约等。

// 原子操作示例：全局求和
__global__ void reduceAtomic(float *g_data, float *g_sum) {
    __shared__ float sdata[256];
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;
    
    sdata[tid] = g_data[idx];
    __syncthreads();
    
    // 块内归约
    for (int s = 1; s < blockDim.x; s *= 2) {
        if (tid % (2 * s) == 0) {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }
    
    // 每个块用原子操作累加
    if (tid == 0) {
        atomicAdd(g_sum, sdata[0]);
    }
}

好了，以上就是CUDA编程基础的核心内容。内存层次让你知道数据该放哪，kernel写法让你知道代码该怎么组织，线程配置决定了并行度，共享内存是性能利器，同步机制保证正确性。这五块内容，是后续学习TensorCore的基石。

想要更多资料，联系可可

📚 课程目录

CUDA编程基础回顾：GPU内存层次结构、Kernel函数编写、线程块与网格配置、共享内存使用、同步机制

1. GPU内存层次结构：从慢到快，从大到小

2. Kernel函数编写：从Hello World到实战

3. 线程块与网格配置：找到那个甜蜜点

4. 共享内存使用：片上缓存的艺术

5. 同步机制：别让线程跑乱了