AMD GPU寄存器入门指南：从基础概念到实战优化

1次阅读

没有评论

共计 1942 个字符，预计需要花费 5 分钟才能阅读完成。

在 AMD GPU 架构中，寄存器（Register）是位于计算单元（CU）内部的超高速存储单元，每个流处理器（SIMD）可直接访问。它们的主要特点包括：

零延迟访问：比全局内存快 100 倍以上
私有性：每个工作项（work-item）拥有独立寄存器组
稀缺资源：总数量受硬件限制（通常每个 CU 仅几千个）

寄存器对性能的影响主要体现在：

减少内存访问瓶颈
提高指令级并行度
降低线程切换开销

标量寄存器（SGPR）：存储地址、循环计数等控制数据
向量寄存器（VGPR）：存储计算涉及的向量数据

静态分配：编译器根据变量声明自动分配
动态分配 ：通过__private 限定符显式控制

32 位宽度（支持 packed data 类型）
支持跨 lane 广播（wavefront 内共享）
占用情况可通过 rocminfo 命令查看

以下 HIP 代码展示寄存器优化前后的矩阵乘法实现对比：

// 基础版本（未优化寄存器使用）__global__ void matmul_naive(float *C, float *A, float *B, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; ++k) {sum += A[row * K + k] * B[k * N + col]; // 反复访问全局内存
        }
        C[row * N + col] = sum;
    }
}

// 优化版本（利用寄存器缓存数据）__global__ void matmul_opt(float *C, float *A, float *B, int M, int N, int K) {__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

    float sum = 0.0f; // 显式使用寄存器
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    for (int k = 0; k < K; k += BLOCK_SIZE) {
        // 将数据块加载到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * K + k + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * N + col];
        __syncthreads();

        // 使用寄存器暂存计算结果
        float reg_cache = 0.0f;
        for (int i = 0; i < BLOCK_SIZE; ++i) {reg_cache += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        }
        sum += reg_cache;
        __syncthreads();}

    if (row < M && col < N) {C[row * N + col] = sum;
    }
}