A100上CUDA核函数L2 Fabric缓存命中率异常问题咨询

我正在使用Nsight Compute对A100上的一个只读CUDA核函数做性能分析。该核函数使用ld.global.cs指令加载全局内存数据，选择.cs是因为处理的是流式数据，无需数据复用。我刻意设置了相邻线程内存访问步长为1KB（大于L2缓存行大小），确保每个内存地址仅被加载一次，因此预期L2缓存命中率为0%。

但实际性能分析结果显示L2缓存命中率为50%，与预期不符。L1/TEX加载请求数符合预期（计算值为3538944 = 216个block × 1024个thread × 16次内存访问），问题似乎出在L2 Fabric上。我想知道：L2 Fabric的额外请求来自何处？已知A100的L2缓存分为两个分区，分区间会有数据传输，但这种场景下的命中率是如何计算的？我认为总请求数应该等于L1的请求数3538944，命中率应该为0%。

核函数代码 #

#include <cstdint>
#include <cuda.h>
#include <cuda_runtime.h>
#include <iostream>

const int BLOCK = 1024;
const int BENCH_SIZE = (1lu << 26);
const int THREAD_STRIDE = (1lu << 16);
const int BLOCK_STRIDE = (1lu << 8);
const int BENCH_ITER = 16;

#define checkCudaErrors(err)  __checkCudaErrors (err, __FILE__, __LINE__)

inline void __checkCudaErrors( CUresult err, const char *file, const int line )
{
    if( CUDA_SUCCESS != err) {
        fprintf(stderr,
                "CUDA Driver API error = %04d from file <%s>, line %i.\n",
                err, file, line );
        exit(-1);
    }
}

__device__ __forceinline__
int ldg_cs_v1(const void *ptr) {
    int ret;
    asm volatile (
        "ld.global.cs.b32 %0, [%1];"
        : "=r"(ret)
        : "l"(ptr)
    );

    return ret;
}

__device__ __forceinline__
void stg_cs_v1(const int &reg, void *ptr) {
    asm volatile (
        "st.global.cs.b32 [%1], %0;"
        : : "r"(reg), "l"(ptr)
    );
}

__global__ void read_kernel(const void *x, void *y) {
    for (int i = 0; i < BENCH_ITER; i++)
    {
        uint32_t idx = BENCH_SIZE * i + blockIdx.x * BLOCK_STRIDE + threadIdx.x * THREAD_STRIDE;
        const int *ldg_ptr = (const int *)x + idx;
        int reg;

        reg = ldg_cs_v1(ldg_ptr);
        // 防止编译器优化掉LDG操作
        if (reg != 0) {
            stg_cs_v1(reg, (int*)y);
        }
    }
}

int main() {
    size_t size_in_byte = (1lu << 30) * 16; // 16GB

    int numBlocks;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, read_kernel, BLOCK, 0);
    printf("Blocknum per SM: %d\n", numBlocks);

    char *ws;
    cudaMalloc(&ws, size_in_byte);
    // 初始化全局内存为0，适配只读核函数逻辑
    cudaMemset(ws, 0, size_in_byte);

    const int L2_FLUSH_SIZE = (1 << 20) * 128;
    int *l2_flush;
    cudaMalloc(&l2_flush, L2_FLUSH_SIZE);
    cudaMemset(l2_flush, 0, L2_FLUSH_SIZE);

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    float time_ms = 0.f;

    cudaEventRecord(start);
    read_kernel<<<numBlocks * 108, BLOCK>>>(ws, nullptr);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time_ms, start, stop);

    printf("time_ms: %f\n", time_ms);
    cudaFree(ws);

    return 0;
}

复现命令 #

nvcc test.cu -arch sm_80 -lcuda -o test
./test
ncu -f -o read_write_kernel --set full --section MemoryWorkloadAnalysis_Chart --replay-mode application ./test

内容的提问来源于stack exchange，提问作者Shulai