1. 理解GPU内存层次与带宽

任务目标

编写一个CUDA内核，分别测量GPU全局内存、共享内存和寄存器的有效带宽与访问延迟，通过对比结果建立对内存层次影响推理性能的直观认识。本案例聚焦于内存访问测量本身，不涉及具体推理网络。

环境依赖

  • NVIDIA GPU（支持CUDA），驱动版本≥525
  • CUDA Toolkit ≥ 11.8（nvcc编译器）
  • Linux或Windows系统，终端执行命令

核心实现

以下代码实现三种内存层次的带宽测量：全局内存通过连续大数组拷贝测峰值带宽；共享内存通过块内数据搬移测带宽；寄存器通过纯寄存器运算测等效带宽。延迟测量采用指针追逐（pointer chasing）模式，分别构造全局内存和共享内存的链式访问。

// memory_bench.cu
#include <cstdio>
#include <cuda_runtime.h>
#include <chrono>

#define CHECK_CUDA(call) do { \
    cudaError_t e = (call); \
    if (e != cudaSuccess) { \
        printf("CUDA error %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(e)); \
        exit(1); \
    } \
} while(0)

// ---------- 带宽测量内核 ----------
__global__ void global_bw_kernel(const float* in, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    for (int i = idx; i < n; i += stride) {
        out[i] = in[i] * 2.0f;  // 读+写
    }
}

__global__ void shared_bw_kernel(const float* in, float* out, int n) {
    extern __shared__ float smem[];
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;
    int stride = gridDim.x * blockDim.x;
    
    for (int i = idx; i < n; i += stride) {
        smem[tid] = in[i];
        __syncthreads();
        out[i] = smem[tid] * 2.0f;
        __syncthreads();
    }
}

__global__ void reg_bw_kernel(const float* in, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    float r0, r1, r2, r3;
    
    for (int i = idx; i < n; i += stride) {
        r0 = in[i];
        r1 = r0 * 2.0f;
        r2 = r1 + 1.0f;
        r3 = r2 * 0.5f;
        out[i] = r3;
    }
}

// ---------- 延迟测量内核（指针追逐） ----------
__global__ void global_latency_kernel(int* ptr, int* result, int iterations) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= 1024) return;  // 限制线程数
    
    int p = idx;
    for (int i = 0; i < iterations; i++) {
        p = ptr[p];  // 依赖链式访问
    }
    result[idx] = p;
}

__global__ void shared_latency_kernel(int* ptr, int* result, int iterations) {
    __shared__ int sptr[1024];
    int tid = threadIdx.x;
    if (tid < 1024) sptr[tid] = ptr[tid];
    __syncthreads();
    
    int p = tid;
    for (int i = 0; i < iterations; i++) {
        p = sptr[p];  // 共享内存链式访问
    }
    result[tid] = p;
}

// ---------- 辅助函数 ----------
float measure_bw(void (*kernel)(const float*, float*, int), 
                 const float* d_in, float* d_out, int n, int threads) {
    int blocks = (n + threads - 1) / threads;
    if (blocks > 65535) blocks = 65535;
    
    cudaEvent_t start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);
    
    // 预热
    kernel<<<blocks, threads>>>(d_in, d_out, n);
    cudaDeviceSynchronize();
    
    cudaEventRecord(start);
    kernel<<<blocks, threads>>>(d_in, d_out, n);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float ms = 0;
    cudaEventElapsedTime(&ms, start, stop);
    
    // 带宽 = 读+写字节数 / 时间
    float bytes = 2.0f * n * sizeof(float);  // 读一次写一次
    float bw = bytes / (ms * 1e-3f) / 1e9f;  // GB/s
    
    cudaEventDestroy(start); cudaEventDestroy(stop);
    return bw;
}

float measure_latency(void (*kernel)(int*, int*, int), 
                      int* d_ptr, int* d_result, int iterations) {
    cudaEvent_t start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);
    
    kernel<<<1, 1024>>>(d_ptr, d_result, iterations);
    cudaDeviceSynchronize();
    
    cudaEventRecord(start);
    kernel<<<1, 1024>>>(d_ptr, d_result, iterations);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float ms = 0;
    cudaEventElapsedTime(&ms, start, stop);
    
    // 延迟 = 总时间 / (迭代次数 * 链长)
    float latency_ns = ms * 1e6f / iterations;
    
    cudaEventDestroy(start); cudaEventDestroy(stop);
    return latency_ns;
}

int main() {
    const int N = 1 << 24;  // 16M floats = 64MB
    const int THREADS = 256;
    const int ITERATIONS = 10000;
    
    // 分配内存
    float *h_in = new float[N];
    float *h_out = new float[N];
    for (int i = 0; i < N; i++) h_in[i] = 1.0f;
    
    float *d_in, *d_out;
    CHECK_CUDA(cudaMalloc(&d_in, N * sizeof(float)));
    CHECK_CUDA(cudaMalloc(&d_out, N * sizeof(float)));
    CHECK_CUDA(cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice));
    
    // 延迟测试用指针数组
    int *h_ptr = new int[1024];
    int *d_ptr, *d_result;
    CHECK_CUDA(cudaMalloc(&d_ptr, 1024 * sizeof(int)));
    CHECK_CUDA(cudaMalloc(&d_result, 1024 * sizeof(int)));
    for (int i = 0; i < 1024; i++) h_ptr[i] = (i + 1) % 1024;  // 循环链
    CHECK_CUDA(cudaMemcpy(d_ptr, h_ptr, 1024 * sizeof(int), cudaMemcpyHostToDevice));
    
    // 测量带宽
    float global_bw = measure_bw(global_bw_kernel, d_in, d_out, N, THREADS);
    
    // 共享内存带宽（使用动态共享内存，每块THREADS个float）
    size_t smem_size = THREADS * sizeof(float);
    cudaFuncSetAttribute(shared_bw_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size);
    float shared_bw = 0;
    {
        cudaEvent_t start, stop;
        cudaEventCreate(&start); cudaEventCreate(&stop);
        int blocks = (N + THREADS - 1) / THREADS;
        if (blocks > 65535) blocks = 65535;
        
        shared_bw_kernel<<<blocks, THREADS, smem_size>>>(d_in, d_out, N);
        cudaDeviceSynchronize();
        
        cudaEventRecord(start);
        shared_bw_kernel<<<blocks, THREADS, smem_size>>>(d_in, d_out, N);
        cudaEventRecord(stop);
        cudaEventSynchronize(stop);
        
        float ms = 0;
        cudaEventElapsedTime(&ms, start, stop);
        shared_bw = 2.0f * N * sizeof(float) / (ms * 1e-3f) / 1e9f;
        cudaEventDestroy(start); cudaEventDestroy(stop);
    }
    
    float reg_bw = measure_bw(reg_bw_kernel, d_in, d_out, N, THREADS);
    
    // 测量延迟
    float global_lat = measure_latency(global_latency_kernel, d_ptr, d_result, ITERATIONS);
    float shared_lat = measure_latency(shared_latency_kernel, d_ptr, d_result, ITERATIONS);
    
    // 输出结果
    printf("=== 内存带宽 (GB/s) ===\n");
    printf("全局内存: %.1f\n", global_bw);
    printf("共享内存: %.1f\n", shared_bw);
    printf("寄存器:   %.1f\n", reg_bw);
    printf("\n=== 访问延迟 (ns) ===\n");
    printf("全局内存: %.1f\n", global_lat);
    printf("共享内存: %.1f\n", shared_lat);
    
    // 清理
    CHECK_CUDA(cudaFree(d_in)); CHECK_CUDA(cudaFree(d_out));
    CHECK_CUDA(cudaFree(d_ptr)); CHECK_CUDA(cudaFree(d_result));
    delete[] h_in; delete[] h_out; delete[] h_ptr;
    
    return 0;
}

编译与运行

nvcc -O3 -arch=sm_80 memory_bench.cu -o memory_bench
./memory_bench

若GPU架构不同（如RTX 30系为sm_86，A100为sm_80），请将-arch参数替换为对应架构编号。

验证方法

程序输出三组带宽和两组延迟数据。判定通过的参考标准：

  • 带宽关系：共享内存带宽应显著高于全局内存（通常高3～10倍）；寄存器带宽通常最高，但受限于指令吞吐，可能仅略高于共享内存
  • 延迟关系：全局内存延迟应在200～600ns范围（取决于GPU型号），共享内存延迟应低于50ns
  • 数值合理性：全局内存带宽应接近GPU理论峰值（如A100约2TB/s，RTX 3090约936GB/s）的60%～85%

若输出结果符合上述数量级关系，说明测量有效，内存层次对性能的影响已得到直观验证。不同GPU型号的具体数值会有差异，但相对关系不变。

常见问题

  • 共享内存带宽偏低：检查是否因__syncthreads()造成过度同步开销，可尝试增大每线程处理的数据量（如每次循环处理4个元素）
  • 寄存器带宽测量不准确：编译器可能将连续运算优化为常量折叠，导致测量值虚高。可尝试使用-O0编译或增加运算复杂度
  • 延迟测量偏差：指针追逐链长度（1024）需大于GPU内存页大小，否则可能命中TLB造成测量值偏小
← 上一章 📖 返回目录 下一章 →