CUDA 内存层次结构

内存优化�?CUDA 性能调优中最重要的一环。理解各级内存的容量、延迟、带宽和使用方式，是写出高性能 Kernel 的关键�?

内存层次总览

线程私有                    Block 共享              全局共享
─────────                  ──────────              ──────────
寄存�?(Register)          共享内存 (Shared Mem)   全局内存 (Global Mem)
本地内存 (Local Mem)                               常量内存 (Constant Mem)
                                                   纹理内存 (Texture Mem)
                                                   L2 Cache

各级内存对比

内存类型	位置	延迟	带宽	容量	生命周期
寄存�?	片上	~1 cycle	极高	255 reg/thread	线程
共享内存	片上	~20 cycles	~19 TB/s/SM	48-164 KB/SM	Block
L1 Cache	片上	~20 cycles	自动	32-128 KB/SM	自动
L2 Cache	片上	~200 cycles	~3 TB/s	40 MB (A100)	自动
全局内存	片外 HBM	~600 cycles	2 TB/s	40-80 GB	应用
常量内存	片外+Cache	~1 cycle(命中)	�?	64 KB	应用
本地内存	片外	~600 cycles	�?	512 KB/thread	线程

寄存器（Register�?

寄存器是最快的存储，每个线程独有�?

__global__ void kernel() {
    // 这些变量存储在寄存器�?
    float a = 1.0f;
    int   b = 42;
    float c = a + b;  // 纯寄存器操作，~1 cycle
}

寄存器溢出（Register Spilling�?

�?Kernel 使用的寄存器超过限制时，溢出�?本地内存*（实际是全局内存，速度极慢）：

# 查看 Kernel 寄存器使用量
nvcc --ptxas-options=-v my_kernel.cu

# 输出示例�?
# ptxas info: Used 32 registers, 1024 bytes smem, 352 bytes cmem[0]

// 限制寄存器使用（可能降低性能，但提高 Occupancy�?
__global__ __launch_bounds__(256, 4)  // maxThreadsPerBlock=256, minBlocksPerSM=4
void kernel() { ... }

// 或编译时指定
// nvcc --maxrregcount=32 my_kernel.cu

共享内存（Shared Memory�?

共享内存是程序员控制的片上缓存，同一 Block 内所有线程共享，速度接近寄存器�?

静态分�?

__global__ void kernel() {
    // 静态共享内存（编译时确定大小）
    __shared__ float smem[256];
    __shared__ int   counter;
    
    smem[threadIdx.x] = globalData[blockIdx.x * blockDim.x + threadIdx.x];
    __syncthreads();
    
    // 现在可以访问 Block 内任意线程写入的数据
    float neighbor = smem[(threadIdx.x + 1) % blockDim.x];
}

动态分�?

// Kernel 定义
__global__ void kernel(int N) {
    extern __shared__ float smem[];  // 大小在启动时指定
    smem[threadIdx.x] = ...;
}

// 启动时指定共享内存大小（字节�?
kernel<<<grid, block, 256 * sizeof(float)>>>(N);

Bank Conflict（存储体冲突�?

共享内存分为 32 �?Bank，每�?Bank 宽度 4 字节。同一 Warp 内多个线程访问同一 Bank 的不同地址时，发生 Bank Conflict，访问被串行化�?

// 无冲突：连续访问，每个线程访问不�?Bank
__shared__ float smem[32];
float val = smem[threadIdx.x];  // 线程 i 访问 Bank i

// 2-way 冲突：步长为 2，线�?0 �?16 都访�?Bank 0
float val = smem[threadIdx.x * 2];

// 32-way 冲突（最坏情况）：所有线程访问同一 Bank
float val = smem[0];  // 广播除外（广播无冲突�?

// 解决方案：添�?padding
__shared__ float smem[32][33];  // 每行多一个元素，错开 Bank
float val = smem[threadIdx.y][threadIdx.x];

共享内存配置

Ampere 架构每个 SM �?192 KB 统一数据缓存，可配置 L1/共享内存比例�?

// 设置共享内存大小（字节）
cudaFuncSetAttribute(myKernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize, 
    163840);  // 160 KB

// 或设�?L1/Shared 比例
cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);  // 偏向共享内存
cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);      // 偏向 L1

全局内存（Global Memory�?

全局内存�?GPU 的主存（HBM），容量最大但延迟最高�?

合并访问（Coalesced Access�?

最重要的全局内存优化原则：同一 Warp 内的线程应访问连续的内存地址，使多次访问合并为一次内存事务�?

// 合并访问（最优）：线�?i 访问地址 base + i
// Warp �?32 个线程访�?128 字节连续内存 �?1 次内存事�?
__global__ void good(float* data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx];  // �?连续访问
}

// 非合并访问（最差）：步长访�?
// 每个线程访问不同 Cache Line �?32 次内存事�?
__global__ void bad(float* data, int stride) {
    int idx = threadIdx.x * stride;
    float val = data[idx];  // �?步长访问
}

// 矩阵转置的访问模式问�?
__global__ void transpose(float* in, float* out, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 读：合并（按行读）✓
    // 写：非合并（按列写）�?
    out[col * N + row] = in[row * N + col];
}

// 优化：用共享内存做中�?
__global__ void transposeFast(float* in, float* out, int N) {
    __shared__ float tile[16][17];  // +1 避免 Bank Conflict
    
    int x = blockIdx.x * 16 + threadIdx.x;
    int y = blockIdx.y * 16 + threadIdx.y;
    
    tile[threadIdx.y][threadIdx.x] = in[y * N + x];  // 合并�?
    __syncthreads();
    
    x = blockIdx.y * 16 + threadIdx.x;
    y = blockIdx.x * 16 + threadIdx.y;
    out[y * N + x] = tile[threadIdx.x][threadIdx.y];  // 合并�?
}

常量内存（Constant Memory�?

常量内存适合存储所有线程都读取的只读数据，有专�?Cache�?

// 声明常量内存（最�?64 KB�?
__constant__ float filter[256];
__constant__ struct Config cfg;

// 从主机写�?
cudaMemcpyToSymbol(filter, h_filter, 256 * sizeof(float));

// Kernel 中读取（广播：所有线程读同一地址时，只需一次内存事务）
__global__ void conv(float* data, float* out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0;
    for (int k = 0; k < 256; k++) {
        sum += data[idx + k] * filter[k];  // filter 从常量缓存读�?
    }
    out[idx] = sum;
}

纹理内存（Texture Memory�?

纹理内存针对 2D 空间局部性优化，适合图像处理�?

// 现代 CUDA 推荐使用 Texture Object API
cudaTextureObject_t texObj;

cudaResourceDesc resDesc = {};
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = d_data;
resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
resDesc.res.linear.desc.x = 32;
resDesc.res.linear.sizeInBytes = N * sizeof(float);

cudaTextureDesc texDesc = {};
texDesc.readMode = cudaReadModeElementType;

cudaCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);

// Kernel 中使�?
__global__ void kernel(cudaTextureObject_t tex, float* out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    out[idx] = tex1Dfetch<float>(tex, idx);
}

cudaDestroyTextureObject(texObj);

内存优化策略总结

优化优先级（从高到低）：

1. 合并全局内存访问
   �?确保 Warp 内线程访问连续地址

2. 最大化共享内存利用
   �?将热点数据放入共享内存，减少全局内存访问

3. 避免 Bank Conflict
   �?共享内存访问�?padding

4. 使用常量内存
   �?只读广播数据放入常量内存

5. 减少寄存器溢�?
   �?控制 Kernel 复杂度，使用 __launch_bounds__

6. 使用异步内存拷贝
   �?cudaMemcpyAsync + Stream 重叠传输与计�?

---

下一篇：cuBLAS �?线性代数库 →