CUDA-内存访问模式

news2025/3/14 10:59:26

在 GPU 计算中，内存访问模式 直接影响程序的性能，尤其是 全局内存（global memory） 访问的合并性（coalescing）和局部性（locality）。

1. GPU 内存层次结构

GPU 具有多级存储，每一级的访问速度不同：

寄存器（Register）：每个线程私有，访问最快。
共享内存（Shared Memory）：线程块（Block）内共享，比全局内存快很多。
L1/L2 缓存（Cache）：缓存局部数据，减少全局内存访问。
全局内存（Global Memory）：最慢，但存储量最大，需要优化访问模式。

2. 主要内存访问模式

(1) 合并访问（Coalesced Access）

✅ 线程按顺序访问连续的内存地址

每个 warp（32 个线程）共同访问一段连续的全局内存。
减少内存事务（memory transactions），提高带宽利用率。

示例（合并访问）：

__global__ void coalesced_access(float *A, float *B, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N)
        B[tid] = A[tid];  // 线程按顺序访问连续地址
}

📌 优化要点：

按行加载（Row-major order） 比按列加载更好。
让线程 ID 与数据索引匹配。

(2) 非合并访问（Non-Coalesced Access）

🚨 线程访问不连续的内存地址

每个线程访问的地址分散在不同的内存块，导致多个内存事务，降低效率。

示例（非合并访问）：

__global__ void non_coalesced_access(float *A, float *B, int N, int stride) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid * stride < N)
        B[tid] = A[tid * stride];  // 线程访问间隔 stride
}

🚩 问题：

内存事务数增加，吞吐量降低。
不建议使用 stride 访问全局内存。

(3) 行主序（Row-Major） vs. 列主序（Column-Major）访问

行主序存储（Row-major order）：C 语言默认方式，数据按行优先存储。
列主序存储（Column-major order）：Fortran、MATLAB 采用，数据按列优先存储。

📌 访问模式建议：

矩阵遍历时，按存储方式优化访问方式，避免列访问全局内存，可用共享内存优化。

示例（行访问更快）：

// 行访问（合并访问）
for (int row = 0; row < N; row++)
    for (int col = 0; col < N; col++)
        B[row * N + col] = A[row * N + col];

// 列访问（非合并访问，性能低）
for (int col = 0; col < N; col++)
    for (int row = 0; row < N; row++)
        B[row * N + col] = A[col * N + row];  // 非合并访问

(4) 共享内存优化

使用 共享内存（shared memory） 解决非合并访问问题：

__global__ void shared_memory_access(float *A, float *B, int N) {
    __shared__ float tile[32][32];  // 共享内存 tile
    int row = threadIdx.y + blockIdx.y * blockDim.y;
    int col = threadIdx.x + blockIdx.x * blockDim.x;

    if (row < N && col < N) {
        tile[threadIdx.y][threadIdx.x] = A[row * N + col];  // 先加载到共享内存
        __syncthreads();
        B[col * N + row] = tile[threadIdx.y][threadIdx.x];  // 从共享内存存回
    }
}

✅ 共享内存缓存局部数据，提高访问效率。

3. 访问模式优化策略

✅ 合并访问：线程访问连续地址，提高全局内存带宽利用率。
✅ 共享内存：减少不规则访问，提高数据重用。
✅ 优化存储顺序：按行遍历 比 按列遍历 快。
✅ 减少 stride 访问：避免跨步访问全局内存。
✅ 使用 L1 缓存（L1 cache）：启用 L1 缓存优化非合并访问。

4. 结论

访问模式	内存层级	访问效率	典型应用
合并访问	全局内存	✅ 高效	标量运算，向量计算
非合并访问	全局内存	❌ 低效	矩阵按列访问
共享内存	共享内存	✅ 高效	矩阵转置、卷积计算
寄存器	寄存器	✅ 最高	线程局部变量