GPU 内存层次结构

GPU 的内存层次结构是其高性能的关键。理解 GPU 的内存层次结构对于编写高效的 CUDA 程序至关重要。本文将详细介绍 GPU 的内存层次结构及其优化策略。

GPU 内存层次概览

GPU 的内存层次结构从快到慢依次为：

内存类型	访问速度	容量	作用域	缓存
Registers（寄存器）	最快	极小（每个线程）	线程私有	无
Shared Memory（共享内存）	极快	小（每个 Block）	Block 共享	无
L1 Cache	快	小（每个 SM）	Block 共享	有
L2 Cache	中等	中（整个 GPU）	全局	有
Global Memory（全局内存）	慢	大（整个 GPU）	全局	有（L1/L2）
Constant Memory（常量内存）	快	小（64KB）	全局	有
Texture Memory（纹理内存）	快	中	全局	有

寄存器（Registers）

寄存器是 GPU 上最快的存储器，每个线程都有自己的一组寄存器。

寄存器特性

• 速度：访问延迟约 1 个时钟周期
• 容量：每个 SM 有数千个寄存器，每个线程分配一部分
• 作用域：线程私有，其他线程无法访问
• 生命周期：与线程生命周期相同

寄存器使用示例

__global__ void kernel(float *data) {
    // 局部变量通常存储在寄存器中
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float temp = data[idx];  // temp 可能存储在寄存器中
    temp *= 2.0f;
    data[idx] = temp;
}

寄存器优化

• 减少寄存器使用可以提高占用率
• 使用 __launch_bounds__ 指定最大线程数
• 避免过多的局部变量

// 指定每个 Block 最多 256 个线程，每个 SM 最多 2 个 Block
__global__ void __launch_bounds__(256, 2) optimizedKernel(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = idx;
}

共享内存（Shared Memory）

共享内存是 SM 内部的高速存储器，同一个 Block 内的线程可以共享访问。

共享内存特性

• 速度：比全局内存快 100 倍以上
• 容量：通常 48KB-96KB（可配置）
• 作用域：Block 内线程共享
• 生命周期：与 Block 生命周期相同

共享内存使用示例

#define TILE_SIZE 16

__global__ void matrixMul(float *A, float *B, float *C, int N) {
    // 声明共享内存
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];

    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;

    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;

    float sum = 0.0f;

    for (int t = 0; t < N / TILE_SIZE; ++t) {
        // 加载数据到共享内存
        tileA[ty][tx] = A[row * N + t * TILE_SIZE + tx];
        tileB[ty][tx] = B[(t * TILE_SIZE + ty) * N + col];

        // 等待所有线程完成加载
        __syncthreads();

        // 使用共享内存中的数据进行计算
        for (int k = 0; k < TILE_SIZE; ++k) {
            sum += tileA[ty][k] * tileB[k][tx];
        }

        // 等待所有线程完成计算
        __syncthreads();
    }

    C[row * N + col] = sum;
}

共享内存优化

• Bank Conflict：避免 Bank 冲突，提高访问效率
• Pad：使用 padding 避免 Bank 冲突
• 异步拷贝：使用 __cp_async 异步加载数据

// 避免 Bank Conflict 的示例
__global__ void avoidBankConflict(float *data) {
    __shared__ float s_data[32];  // 32 个 Bank，每个 Bank 4 字节

    int tid = threadIdx.x;
    s_data[tid] = data[tid];  // 每个线程访问不同 Bank，无冲突

    __syncthreads();

    data[tid] = s_data[tid];
}

// 使用 Padding 避免 Bank Conflict
__global__ void withPadding(float *data) {
    __shared__ float s_data[32][33];  // 添加 Padding

    int tid = threadIdx.x;
    s_data[tid][0] = data[tid];  // 避免 Bank 冲突

    __syncthreads();

    data[tid] = s_data[tid][0];
}

全局内存（Global Memory）

全局内存是 GPU 上最大的存储空间，但访问速度最慢。

全局内存特性

• 速度：访问延迟高（数百个时钟周期）
• 容量：大（几 GB 到几十 GB）
• 作用域：全局，所有线程都可以访问
• 缓存：有 L1/L2 缓存

内存合并访问

当相邻线程访问连续的内存地址时，GPU 可以将多个内存请求合并为一个事务。

// 好的内存合并
__global__ void coalescedAccess(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = idx;  // 连续访问，合并
}

// 不好的内存访问
__global__ void stridedAccess(float *data, int stride) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx * stride] = idx;  // 跨步访问，不合并
}

全局内存优化

• 使用行优先存储（Row-major）
• 避免跨步访问（Strided Access）
• 使用共享内存减少全局内存访问
• 考虑数据对齐（Alignment）

常量内存（Constant Memory）

常量内存是只读内存，有缓存优化，适合存储常量数据。

常量内存特性

• 速度：快，有缓存
• 容量：64KB
• 作用域：全局
• 访问：只读

常量内存使用示例

// 声明常量内存
__constant__ float constData[256];

// 从主机复制数据到常量内存
cudaMemcpyToSymbol(constData, hostData, 256 * sizeof(float));

// 在 Kernel 中使用
__global__ void useConstant(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = constData[idx % 256];  // 从常量内存读取
}

纹理内存（Texture Memory）

纹理内存是专门为图形渲染优化的内存，也适用于某些计算场景。

纹理内存特性

• 速度：快，有缓存
• 容量：中
• 作用域：全局
• 特性：支持 2D/3D 寻址、插值、边界处理

纹理内存使用示例

// 声明纹理引用
texture texRef;

// 绑定纹理到全局内存
cudaBindTexture(NULL, texRef, d_data, width * height * sizeof(float));

// 在 Kernel 中使用纹理
__global__ void useTexture(float *output) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    // 使用纹理坐标读取
    float u = x / (float)width;
    float v = y / (float)height;
    output[y * width + x] = tex2D(texRef, u, v);  // 纹理读取
}

// 解绑纹理
cudaUnbindTexture(texRef);

内存访问模式优化

1. 内存合并

确保相邻线程访问连续的内存地址。

2. 避免 Bank 冲突

在共享内存中，避免多个线程同时访问同一个 Bank。

3. 使用共享内存

利用共享内存减少全局内存访问次数。

4. 数据重用

将频繁使用的数据加载到共享内存或寄存器中。

内存管理函数

// 分配全局内存
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));

// 释放全局内存
cudaFree(d_data);

// 内存拷贝：主机到设备
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// 内存拷贝：设备到主机
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);

// 内存拷贝：设备到设备
cudaMemcpy(d_dest, d_src, N * sizeof(float), cudaMemcpyDeviceToDevice);

// 内存填充
cudaMemset(d_data, 0, N * sizeof(float));

// 异步内存拷贝
cudaMemcpyAsync(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice, stream);

总结

GPU 的内存层次结构是其高性能的关键。理解不同内存类型的特性和访问模式，可以帮助我们编写高效的 CUDA 程序：

• 寄存器：最快的存储，用于局部变量
• 共享内存：Block 内共享，用于数据重用
• 全局内存：最大的存储，但访问慢，需要优化访问模式
• 常量/纹理内存：只读内存，有缓存优化

通过合理使用不同层次的内存，可以充分发挥 GPU 的计算能力，实现高性能的并行计算。