迈向可编程观测：在GPU Kernel中构建类eBPF风格的性能探针

引言

第一部分：CUDA架构基础，是理解性能瓶颈的起点

#include<iostream>#include<vector>#include<random>#include<stdexcept>#include<cmath>#include<cuda_runtime.h>
#define CUDA_CHECK(err) { \    cudaError_t err_ = (err); \    if (err_ != cudaSuccess) { \        std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \        std::cerr << ": " << cudaGetErrorString(err_) << std::endl; \        exit(EXIT_FAILURE); \    } \}
voidgemm_cpu(constfloat* A, constfloat* B, float* C, int M, int N, int K){    for (int r = 0; r < M; ++r) {        for (int c = 0; c < N; ++c) {            float sum = 0.0f;            for (int i = 0; i < K; ++i) {                sum += A[r * K + i] * B[i * N + c];            }            C[r * N + c] = sum;        }    }}
voidverify_result(constfloat* C_gpu, constfloat* C_cpu, int M, int N){    double max_error = 0.0;    for (int i = 0; i < M * N; ++i) {        max_error = fmax(max_error, fabs(C_gpu[i] - C_cpu[i]));    }    std::cout << "Max error between GPU and CPU results: " << max_error << std::endl;    if (max_error > 1e-4) {        std::cout << "Verification FAILED!" << std::endl;    } else {        std::cout << "Verification PASSED!" << std::endl;    }}
__global__ voidgemm_gpu(constfloat* A, constfloat* B, float* C, 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 i = 0; i < K; ++i) {            sum += A[row * K + i] * B[i * N + col];        }        C[row * N + col] = sum;    }}
intmain(int argc, char** argv){    int M = 1024;    int N = 1024;    int K = 1024;
    std::cout << "Matrix dimensions: C(" << M << "x" << N << ") = A(" << M << "x" << K << ") * B(" << K << "x" << N << ")" << std::endl;
    std::vector<float> h_A(M * K);    std::vector<float> h_B(K * N);    std::vector<float> h_C_gpu(M * N, 0);      std::vector<float> h_C_cpu(M * N, 0); 
    std::random_device rd;    std::mt19937 gen(rd());    std::uniform_real_distribution<> dis(0.0, 1.0);    for (int i = 0; i < M * K; ++i) h_A[i] = dis(gen);    for (int i = 0; i < K * N; ++i) h_B[i] = dis(gen);
    float *d_A, *d_B, *d_C;    CUDA_CHECK(cudaMalloc(&d_A, M * K * sizeof(float)));    CUDA_CHECK(cudaMalloc(&d_B, K * N * sizeof(float)));    CUDA_CHECK(cudaMalloc(&d_C, M * N * sizeof(float)));    CUDA_CHECK(cudaMemcpy(d_A, h_A.data(), M * K * sizeof(float), cudaMemcpyHostToDevice));    CUDA_CHECK(cudaMemcpy(d_B, h_B.data(), K * N * sizeof(float), cudaMemcpyHostToDevice));
    int recommend_blocksize;     int recommend_gridsize;    cudaOccupancyMaxPotentialBlockSize(        &recommend_gridsize,        &recommend_blocksize,        (void*)gemm_gpu,        0,        0);    int device;    cudaDeviceProp prop;    cudaGetDevice(&device);    cudaGetDeviceProperties(&prop, device);    int dimX = prop.warpSize;    int dimY = 1;    if (recommend_blocksize % dimX == 0){        dimY = recommend_blocksize/dimX;    }    else{        dimX = static_cast<int>(sqrt(recommend_blocksize));        dimY = static_cast<int>(sqrt(recommend_blocksize));    }    dim3 block_dim(dimX, dimY);    dim3 grid_dim((N + block_dim.x - 1) / block_dim.x, (M + block_dim.y - 1) / block_dim.y);
    std::cout << "\n--- Running GEMM ---" << std::endl;    gemm_gpu<<<grid_dim, block_dim>>>(d_A, d_B, d_C, M, N, K);    CUDA_CHECK(cudaDeviceSynchronize());    CUDA_CHECK(cudaMemcpy(h_C_gpu.data(), d_C, M * N * sizeof(float), cudaMemcpyDeviceToHost));    gemm_cpu(h_A.data(), h_B.data(), h_C_cpu.data(), M, N, K);    verify_result(h_C_gpu.data(), h_C_cpu.data(), M, N);
    CUDA_CHECK(cudaFree(d_A));    CUDA_CHECK(cudaFree(d_B));    CUDA_CHECK(cudaFree(d_C));
    return0;}

开始│├─ line55-57   --定义矩阵大小 (M, N, K)│├─ line61-70   --主机内存分配与初始化│├─ line72-75   --GPU 内存分配 (d_A, d_B, d_C)│├─ line76-77   --数据从主机传送到设备 (H→D)│├─ line79-101   --评估最佳block_size，配置block_dim、grid_dim│├─ line104     --启动Kernel gemm_gpu <<<grid, block>>>  │├─ line42-52   --GPU上Kernel执行│├─ line105      --GPU同步等待完成│├─ line106    --结果从设备传回主机 (D→H) → h_C_gpu│├─ line107、17-27     --CPU计算参考 gemm_cpu()│├─ line108、29-40     --对比CPU、GPU计算结果 verify_result()│├─ line110-112  --释放GPU内存│程序结束

内存地址 ->[ a00, a01, a02,  a10, a11, a12 ]^---------------^  ^----------- ^   第0行 (3个元素)    第1行 (3个元素) 

        threadIdx.x->         0   1   2   3   ... 15       +---+---+---+---+---+  0    | 0 | 1 | 2 | 3 |...| 15  ← 第0行 → 线性ID: 0~15  → warp 0ty     +---+---+---+---+---+  1    |16 |17 |18 |19 |...| 31  ← 第1行 → 线性ID: 16~31 → warp 0       +---+---+---+---+---+  2    |32 |33 |...        |     ← 第2行 → 线性ID: 32~47 → warp 1       +---+---+---+---+---+  3    |48 |...            |     ← 第3行 → 线性ID: 48~63 → warp 1       +---+---+---+---+---+  4    |64 |...            |     ← 第4行 → 线性ID: 64~79 → warp 2       +---+---+---+---+---+...

#include<iostream>#include<vector>#include<random>#include<stdexcept>#include<cmath>#include<cuda_runtime.h>#define CUDA_CHECK(err) { \    cudaError_t err_ = (err); \    if (err_ != cudaSuccess) { \        std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \        std::cerr << ": " << cudaGetErrorString(err_) << std::endl; \        exit(EXIT_FAILURE); \    } \}#define TILE_SIZE 32__global__ voidgemm_gpu(constfloat* A, constfloat* B, float* C, int M, int N, int K){    __shared__ float sA[TILE_SIZE][TILE_SIZE];    __shared__ float sB[TILE_SIZE][TILE_SIZE];    int row = blockIdx.y * TILE_SIZE + threadIdx.y;    int col = blockIdx.x * TILE_SIZE + threadIdx.x;    float sum = 0.0;    for (int i = 0; i < (K+TILE_SIZE-1)/TILE_SIZE; ++i)    {        if (row <M && i*TILE_SIZE+threadIdx.x <K)            sA[threadIdx.y][threadIdx.x]=A[row*K+i*TILE_SIZE+threadIdx.x];        else            sA[threadIdx.y][threadIdx.x]=0.0;        if ( i*TILE_SIZE+threadIdx.y <K && col <N)            sB[threadIdx.y][threadIdx.x]=B[(i*TILE_SIZE+threadIdx.y)*N+col];        else            sB[threadIdx.y][threadIdx.x]=0.0;        __syncthreads();        for (int k =0; k< TILE_SIZE; ++k)        {            sum += sA[threadIdx.y][k]*sB[k][threadIdx.x];        }        __syncthreads();    }    if (row < M && col < N) {        C[row * N + col] = sum;    }}voidverify_result(constfloat* C_gpu, constfloat* C_cpu, int M, int N){    double max_error = 0.0;    for (int i = 0; i < M * N; ++i) {        max_error = fmax(max_error, fabs(C_gpu[i] - C_cpu[i]));    }    std::cout << "Max error between GPU and CPU results: " << max_error << std::endl;    if (max_error > 1e-4) {        std::cout << "Verification FAILED!" << std::endl;    } else {        std::cout << "Verification PASSED!" << std::endl;    }}voidgemm_cpu(constfloat* A, constfloat* B, float* C, int M, int N, int K){    for (int r = 0; r < M; ++r) {        for (int c = 0; c < N; ++c) {            float sum = 0.0f;            for (int i = 0; i < K; ++i) {                sum += A[r * K + i] * B[i * N + c];            }            C[r * N + c] = sum;        }    }}intmain(int argc, char** argv){    int M = 1024;    int N = 1024;    int K = 1024;    std::cout << "Matrix dimensions: C(" << M << "x" << N << ") = A(" << M << "x" << K << ") * B(" << K << "x" << N << ")" << std::endl;    std::vector<float> h_A(M * K);    std::vector<float> h_B(K * N);    std::vector<float> h_C_gpu(M * N, 0);      std::vector<float> h_C_cpu(M * N, 0);     std::random_device rd;    std::mt19937 gen(rd());    std::uniform_real_distribution<> dis(0.0, 1.0);    for (int i = 0; i < M * K; ++i) h_A[i] = dis(gen);    for (int i = 0; i < K * N; ++i) h_B[i] = dis(gen);    float *d_A, *d_B, *d_C;    CUDA_CHECK(cudaMalloc(&d_A, M * K * sizeof(float)));    CUDA_CHECK(cudaMalloc(&d_B, K * N * sizeof(float)));    CUDA_CHECK(cudaMalloc(&d_C, M * N * sizeof(float)));    CUDA_CHECK(cudaMemcpy(d_A, h_A.data(), M * K * sizeof(float), cudaMemcpyHostToDevice));    CUDA_CHECK(cudaMemcpy(d_B, h_B.data(), K * N * sizeof(float), cudaMemcpyHostToDevice));    dim3 block_dim(TILE_SIZE, TILE_SIZE);    dim3 grid_dim((N + block_dim.x - 1) / block_dim.x, (M + block_dim.y - 1) / block_dim.y);    std::cout << "\n--- Running GEMM ---" << std::endl;    gemm_gpu<<<grid_dim, block_dim>>>(d_A, d_B, d_C, M, N, K);    CUDA_CHECK(cudaDeviceSynchronize());    CUDA_CHECK(cudaMemcpy(h_C_gpu.data(), d_C, M * N * sizeof(float), cudaMemcpyDeviceToHost));    gemm_cpu(h_A.data(), h_B.data(), h_C_cpu.data(), M, N, K);    verify_result(h_C_gpu.data(), h_C_cpu.data(), M, N);    CUDA_CHECK(cudaFree(d_A));    CUDA_CHECK(cudaFree(d_B));    CUDA_CHECK(cudaFree(d_C));    return0;}

==优化前==---------------------------------------------- ----------- -------------Metric Name                                    Metric Unit  Metric Value---------------------------------------------- ----------- -------------l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum          byte 5,368,709,120l1tex__t_bytes_pipe_lsu_mem_global_op_st.sum          byte     4,194,304l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum      sector   167,772,160l1tex__t_sectors_pipe_lsu_mem_global_op_st.sum      sector       131,072---------------------------------------------- ----------- -------------==优化后==---------------------------------------------- ----------- ------------Metric Name                                    Metric Unit Metric Value---------------------------------------------- ----------- ------------l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum          byte  268,435,456l1tex__t_bytes_pipe_lsu_mem_global_op_st.sum          byte    4,194,304l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum      sector    8,388,608l1tex__t_sectors_pipe_lsu_mem_global_op_st.sum      sector      131,072---------------------------------------------- ----------- ------------

==优化前==Total Blocks: 1376, Total Tasks: 33024Average Running Time: 190207 (cycles)Average Idle Time: 1196 (cycles)==优化后==Total Blocks: 1024, Total Tasks: 32768Average Running Time: 106999 (cycles)Average Idle Time: 474 (cycles)

==TILE_SIZE 32==-------------------------------------------------------- ----------- ------------Metric Name                                              Metric Unit Metric Value-------------------------------------------------------- ----------- ------------l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.avg                 8,733.29l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.max                    9,932l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.min                    8,103l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                  628,797-------------------------------------------------------- ----------- ------------==TILE_SIZE 16==-------------------------------------------------------- ----------- ------------Metric Name                                              Metric Unit Metric Value-------------------------------------------------------- ----------- ------------l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.avg                 2,309.79l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.max                    2,745l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.min                    1,774l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                  166,305-------------------------------------------------------- ----------- ------------==cuBlAS==-------------------------------------------------------- ----------- ------------Metric Name                                              Metric Unit Metric Value-------------------------------------------------------- ----------- ------------l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.avg                 2,275.56l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.max                    2,304l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.min                    2,048l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                  163,840-------------------------------------------------------- ----------- ------------

第二部分：从CUDA到PTX，看懂编译器生成的“中间语言”

.version 8.7                      .target sm_52.address_size 64                    --> 全局声明，标识这段代码的PTX版本，目标架构与地址模式
.visible .entry _Z8gemm_gpuPKfS0_Pfiii(          --> 定义一个可被CPU调用的Kernel函数        .param .u64 _Z8gemm_gpuPKfS0_Pfiii_param_0,     --> 参数声明 A        .param .u64 _Z8gemm_gpuPKfS0_Pfiii_param_1,     --> 参数声明 B        .param .u64 _Z8gemm_gpuPKfS0_Pfiii_param_2,     --> 参数声明 C        .param .u32 _Z8gemm_gpuPKfS0_Pfiii_param_3,     --> 参数声明 M        .param .u32 _Z8gemm_gpuPKfS0_Pfiii_param_4,     --> 参数声明 N        .param .u32 _Z8gemm_gpuPKfS0_Pfiii_param_5      --> 参数声明 K){        .reg .pred      %p<9>;              --> 寄存器声明9个bool寄存器         .reg .f32       %f<30>;              --> 寄存器声明30个float寄存器         .reg .b32       %r<32>;              --> 寄存器声明32个通用32位寄存器        .reg .b64       %rd<34>;            --> 寄存器声明34个通过64位寄存器

        ld.param.u64    %rd18, [_Z8gemm_gpuPKfS0_Pfiii_param_0];  --> 参数加载 A->%rd18        ld.param.u64    %rd19, [_Z8gemm_gpuPKfS0_Pfiii_param_1];  --> 参数加载 B->%rd19        ld.param.u64    %rd17, [_Z8gemm_gpuPKfS0_Pfiii_param_2];  --> 参数加载 C->%rd17        ld.param.u32    %r14, [_Z8gemm_gpuPKfS0_Pfiii_param_3];    --> 参数加载 M->%rd14        ld.param.u32    %r12, [_Z8gemm_gpuPKfS0_Pfiii_param_4];    --> 参数加载 N->%rd12        ld.param.u32    %r13, [_Z8gemm_gpuPKfS0_Pfiii_param_5];    --> 参数加载 K->%rd13        cvta.to.global.u64      %rd1, %rd19;  --> 将通用地址转换为全局内存地址空间的64位指针得B基地址B[0][0]        cvta.to.global.u64      %rd2, %rd18;        mov.u32         %r15, %ntid.y;       --> %ntid.y=blockDim.y,传入%r15        mov.u32         %r16, %ctaid.y;      --> %ctaid.y=blockIdx.y,传入%r16        mov.u32         %r17, %tid.y;      --> %tid.y=threadIdx.y,传入%r17        mad.lo.s32      %r1, %r16, %r15, %r17;  -->  %r16*%r15+%r17=row,计算结果传入%r1        mov.u32         %r18, %ntid.x;              mov.u32         %r19, %ctaid.x;              mov.u32         %r20, %tid.x;              mad.lo.s32      %r2, %r19, %r18, %r20;          setp.ge.s32     %p1, %r1, %r14;      -->  此处用的是%r1是否>=%r14来判断row<M是否成立，判断结果放入%p1        setp.ge.s32     %p2, %r2, %r12;              or.pred         %p3, %p1, %p2;      --> 编译器优化，判断(row>=M||col>=N)        @%p3 bra        $L__BB0_9;        --> 判断%p3结果是否True，如果真跳转到$L__BB0_9(就是ret结束）
        setp.lt.s32     %p4, %r13, 1;        --> 编译器优化，用%r13是否<1判断K是否=0,判断K是否=0        mov.f32         %f29, 0f00000000;    --> 根据代码上下文可知，%f29=sum，赋值0.0f        @%p4 bra        $L__BB0_8;        --> 如果K=0goto $L__BB0_8，不会进入循环体
        add.s32         %r22, %r13, -1;      --> K-1的结果传入%r22        and.b32         %r31, %r13, 3;      --> 按位与计算 K&3，等价于K%4（4路展开是由编译器控制，表示一次循环控制每次迭代处理4次原循环的工作，代码里可以通过#pragma unroll独立指定        setp.lt.u32     %p5, %r22, 3;      --> 判断%r22是否<3,等价于判断K是否小于4        mov.f32         %f29, 0f00000000;        mov.u32         %r30, 0;        --> %r30是i，这里就是给i赋值0        @%p5 bra        $L__BB0_5;        --> 如果K<4为True，goto $L__BB0_5
        sub.s32         %r29, %r13, %r31;    --> （4路展开初始化部分）K-（K%4）传入%r29，这里是判断K是否可以被4整除，用于后续不能被4整除的部分goto  $L__BB0_5。 计算得到的%r29是能被4整除部分，用于后续控制4路循环循环次数        mul.lo.s32      %r24, %r13, %r1;    --> K*row 计算 A[row][i] 的行偏移，传入%r24        mul.wide.s32    %rd3, %r24, 4;      --> 计算A[row][0]的字节偏移        mul.wide.s32    %rd20, %r2, 4;      -->  以行首为基准，计算第col列的字节偏移，即计算B[0][col]的字节偏移        add.s64         %rd30, %rd1, %rd20;     -->  基地址+字节偏移=B[0][col]的起始地址        mul.wide.s32    %rd5, %r12, 4;      --> N*4=B矩阵的每行字节数（stride）        mov.f32         %f29, 0f00000000;        mov.u32         %r30, 0;        mov.u64         %rd31, %rd2;      --> A基地址传入%rd31
$L__BB0_4:                    --> K > 4, 这里是示例代码配置下实际会执行的循环体部分        add.s64         %rd21, %rd31, %rd3;      --> 基地址+字节偏移=A[row][0]起始地址        ld.global.f32   %f12, [%rd30];        -->  从全局内存地址%rd30处读取一个float，放入寄存器%f12        ld.global.f32   %f13, [%rd21];                  fma.rn.f32      %f14, %f13, %f12, %f29;     -->  %f13*%f12+%f29，对于首次循环就是计算sum=A[row][0]*B[0][col]+sum        add.s64         %rd22, %rd30, %rd5;      --> B[1][col]        ld.global.f32   %f15, [%rd22];                  ld.global.f32   %f16, [%rd21+4];      --> A[row][1]        fma.rn.f32      %f17, %f16, %f15, %f14;        add.s64         %rd23, %rd22, %rd5;      --> B[2][col]        ld.global.f32   %f18, [%rd23];        ld.global.f32   %f19, [%rd21+8];      --> A[row][2]        fma.rn.f32      %f20, %f19, %f18, %f17;        add.s64         %rd24, %rd23, %rd5;        add.s64         %rd30, %rd24, %rd5;        ld.global.f32   %f21, [%rd24];        ld.global.f32   %f22, [%rd21+12];        fma.rn.f32      %f29, %f22, %f21, %f20;        add.s32         %r30, %r30, 4;        --> 4路展开所以下一循环从i+4开始        add.s64         %rd31, %rd31, 16;        add.s32         %r29, %r29, -4;        --> 循环次数减4        setp.ne.s32     %p6, %r29, 0;        --> 判断%r29是否减到0了        @%p6 bra        $L__BB0_4;          --> 如果不等于0，goto $L__BB0_4 开始为止继续计算
$L__BB0_5:                    --> K<4 ，普通单路循环，处理掉不够4次的情况        setp.eq.s32     %p7, %r31, 0;              @%p7 bra        $L__BB0_8;
        mad.lo.s32      %r25, %r30, %r12, %r2;          mul.wide.s32    %rd25, %r25, 4;              add.s64         %rd33, %rd1, %rd25;        mul.wide.s32    %rd11, %r12, 4;        mad.lo.s32      %r26, %r13, %r1, %r30;        mul.wide.s32    %rd26, %r26, 4;        add.s64         %rd32, %rd2, %rd26;
$L__BB0_7:        .pragma "nounroll";        ld.global.f32   %f23, [%rd33];        ld.global.f32   %f24, [%rd32];        fma.rn.f32      %f29, %f24, %f23, %f29;        add.s64         %rd33, %rd33, %rd11;        add.s64         %rd32, %rd32, 4;        add.s32         %r31, %r31, -1;        setp.ne.s32     %p8, %r31, 0;        @%p8 bra        $L__BB0_7;
$L__BB0_8:                     --> 对应到是 C[row * N + col] = sum        mad.lo.s32      %r27, %r1, %r12, %r2;  --> 对应 row * N + col        cvta.to.global.u64      %rd27, %rd17;          mul.wide.s32    %rd28, %r27, 4;      --> %r27*4字节，传入%rd28，这是在计算C元素的内存地址        add.s64         %rd29, %rd27, %rd28;          st.global.f32   [%rd29], %f29;      -->  写入全局内存
$L__BB0_9:        ret;
}

第三部分：深入 PTX，构建可编程的GPU性能探针

from neutrino import probe, Mapimport neutrino.language as nl
CALLBACK = "block_sched_callback.py"# for trace analysis
# declare maps for persistence@Map(level="warp", type="array", size=16, cap=1)class block_sched:    start: nl.u64    elapsed: nl.u32    cuid: nl.u32
# declare probe registers shared across probesstart: nl.u64 = 0# starting clockelapsed: nl.u64 = 0# elapsed time, initialized to 0
# define probes with decorator@probe(pos="kernel", level="warp", before=True)def thread_start():    start = nl.clock()
@probe(pos="kernel", level="warp")def thread_end():    elapsed = nl.clock() - start    block_sched.save(start, elapsed, nl.cuid())

./trace/Aug14_112435_2682├── event.log          --> 工具的事件日志├── kernel            --> 目标程序里每个Kernel会生成一个子目录│   ├── 0_d02df43fa8dfe9acebcd959bee6f98c2404e40f9│   │   ├── kernel.info      --> 记录Kernel名称、形状│   │   ├── original.bin    --> 原Kernel的cubin文件│   │   ├── original.ptx    --> 原Kernel的ptx汇编文件│   │   ├── probed.bin      --> 插桩后生成的cubin文件      │   │   ├── probed.ptx      --> 插桩后的ptx汇编文件│   │   ├── process.log      --> ptxas对插桩后的ptx生成cubin的日志记录│   │   ├── pruned.bin│   │   └── pruned.ptx      --> 经过裁剪的仅包含与探针相关逻辑的PTX代码片段│   └── 1_f3396b6dedff3f84b1c26dccb7524fb932dd6c14│       ├── kernel.info│       ├── original.bin│       ├── original.ptx│       ├── probed.bin│       ├── probed.ptx│       ├── process.log│       ├── pruned.bin│       └── pruned.ptx├── probe.toml          --> toml格式的探针文件├── read.py            --> Neutrino预置的用来解析result目录中二进制数据的脚本└── result    ├── 0.297500.bin      --> 通过Map收集到的所有原始二进制数据,每个Kernel一个    └── 5.111352.bin

importstruct
# i表示int32，4字节fmt_header ='iiiiiiii'header_size = struct.calcsize(fmt_header)
fmt_section = 'iii'section_size =  struct.calcsize(fmt_section)
# q表示int64，8字节fmt_data = 'qii'data_size = struct.calcsize(fmt_data)
events = []withopen("0.297500.bin", "rb") as f:    #读取&丢弃    header_data = f.read(header_size)    #读取&丢弃    section_data = f.read(section_size)    offset = struct.unpack(fmt_section, section_data)[2]    #跳转到自定义Map数据起始位置    f.seek(offset)    while chunk := f.read(data_size):        iflen(chunk) == data_size:            start, elapsed, cuid = struct.unpack(fmt_data, chunk)            events.append((start, elapsed, cuid))            print(f"start={start}, elapsed={elapsed}, cuid={cuid}")

# Neutrino Auto-Generated Code for Trace ReadingimportstructfromtypingimportNamedTuple, List, Tuple, DictfromneutrinoimportTraceHeader, TraceSectionclassblock_sched(NamedTuple):    start: int    elapsed: int    cuid: intdef parse(path: str) -> Tuple[TraceHeader, List[TraceSection], Dict[str, List[List[NamedTuple]]]]:    with open(path, "rb") as f:        header: TraceHeader = TraceHeader(*struct.unpack("iiiiiiii", f.read(32)))        sections: List[TraceSection] = []        for _ in range(header.numProbes):            sections.append(TraceSection(*struct.unpack("IIQ", f.read(16))))        gridSize = header.gridDimX * header.gridDimY * header.gridDimZ        blockSize = header.blockDimX * header.blockDimY * header.blockDimZ        records: Dict[str, List[List[NamedTuple]]] = dict()        # Read block_sched        records["block_sched"] = []        f.seek(sections[0].offset)        for i in range(gridSize):            records["block_sched"].append([])            for j in range(blockSize // sections[0].warpDiv):                records["block_sched"][-1].append([])                for k in range(sections[0].size // 16):                    records["block_sched"][i][j].append(block_sched(*struct.unpack("qII", f.read(16))))    return header, sections, records# END of Neutrino Auto-Generated Code for Trace Readingimport sysimport numpy as npfrom typing import NamedTuple, Listheader, sections, records_map = parse(sys.argv[1]) # filled by path to tracerecords = records_map["block_sched"]all_tasks = [task for block in records for warp in block for task in warp]all_tasks.sort(key=lambda task: task.start)unique_sms = sorted(list(set(task.cuid for task in all_tasks)))sm_to_idx = {sm_id: i for i, sm_id in enumerate(unique_sms)}num_sms = len(unique_sms)sm_timelines = [[] for _ in range(num_sms)] sched_times = [0.0] * num_smswork_times = [0.0] * num_smsfor task in all_tasks:    sm_idx = sm_to_idx[task.cuid]    timeline = sm_timelines[sm_idx]    #按照SM处理器粒度，统计线程耗时总和    work_times[sm_idx] += task.elapsed    ifnot timeline:        timeline.append(task)        continue    best_slot = None    earliest_end_time = float('inf')    for old_task in timeline:        if old_task.start + old_task.elapsed <= task.start:            if old_task.start + old_task.elapsed < earliest_end_time:                earliest_end_time = old_task.start + old_task.elapsed                best_slot = old_task    if best_slot is not None:        sched_delay = task.start - (best_slot.start + best_slot.elapsed)        #按照SM处理器粒度，统计Warp空闲时间        sched_times[sm_idx] += sched_delay        timeline.remove(best_slot)        timeline.append(task)    else:        timeline.append(task)avg_work_time = np.array(work_times).sum() / len(all_tasks) if all_tasks else0avg_idle_gap = np.array(sched_times).sum() / len(all_tasks) if all_tasks else0print(f"Total Blocks: {len(records)}, Total Tasks: {len(all_tasks)}")print(f"Average Running Time: {int(avg_work_time)} (cycles)")print(f"Average Idle Time: {int(avg_idle_gap)} (cycles)")

regs = 25CALLBACK = "warp_duration_callback.py"
[probe.warp_duration_trace]level = "warp"pos = "kernel"before = """.reg .b32   %tmp<21>;   .reg .pred  %lead;       mov.u32     %tmp2, %laneid;  setp.eq.u32 %lead, %tmp2, 0; @%lead mov.u32 %tmp3, %nwarpid; @%lead mov.u32 %tmp4, %tid.x;         // threadIdx.x@%lead mov.u32 %tmp5, %tid.y;         // threadIdx.y@%lead mov.u32 %tmp6, %tid.z;         // threadIdx.z@%lead mov.u32 %tmp7, %ntid.x;        // blockDim.x@%lead mov.u32 %tmp8, %ntid.y;        // blockDim.y@%lead mov.u32 %tmp18, %ntid.z;       // blockDim.z@%lead mov.u32 %tmp9, %ctaid.x;       // blockIdx.x@%lead mov.u32 %tmp10, %ctaid.y;      // blockIdx.y@%lead mov.u32 %tmp11, %ctaid.z;      // blockIdx.z@%lead mov.u32 %tmp12, %nctaid.x;     // gridDim.x@%lead mov.u32 %tmp13, %nctaid.y;     // gridDim.y@%lead mad.lo.s32 %tmp14, %tmp8,  %tmp6,  %tmp5;   // blockDim.y * threadIdx.z + threadIdx.y@%lead mad.lo.s32 %tmp15, %tmp14, %tmp7,  %tmp4;   // thread_idx in per block = (blockDim.y * threadIdx.z + threadIdx.y) * blockDim.x + threadIdx.x@%lead div.s32 %tmp15, %tmp15, 32;                 // get persistent warpid instead of dynamic %warpid@%lead mad.lo.s32 %tmp16, %tmp13, %tmp11, %tmp10;  // gridDim.y * blockIdx.z + blockIdx.y@%lead mad.lo.s32 %tmp17, %tmp16, %tmp12, %tmp9;   // block_idx = (gridDim.y * blockIdx.z + blockIdx.y) * gridDim.x + blockIdx.x@%lead mul.lo.s32 %tmp19, %tmp7, %tmp8;@%lead mul.lo.s32 %tmp20, %tmp19, %tmp18;          // thread_num=blockDim.x*blockDim.y*blockDim.z@%lead div.s32 %tmp20, %tmp20, 32;                 // warp_num per block = thread_num/32@%lead mad.lo.s32 %tmp1,  %tmp17, %tmp20,  %tmp15; // buf_idx = block_idx * warp_num + warpIdx@%lead mov.u64 %NR0, %clock64;"""after = """mov.u64 %NR2, %clock64;sub.u64 %NR1, %NR2, %NR0;cvt.u64.u32 %NR3, %tmp1;SAVE [ warp_duration ] { %NR0, %NR1, %NR3 }"""
[map.warp_duration]level = "warp"type = "array"size = 32cap = 1
[map.warp_duration.regs]start = [ "u64", "None",]elapsed = [ "u64", "None",]warpid = [ "u64", "None",]

importstructfromtypingimportNamedTuplefromneutrinoimportTraceHeader, TraceSection

classwarp_duration(NamedTuple):    start: int    elapsed: int    warpid: int

def parse(path: str) -> tuple[TraceHeader, list[TraceSection], dict[str, list[list[NamedTuple]]]]:    with open(path, "rb") as f:        header: TraceHeader = TraceHeader(*struct.unpack("iiiiiiii", f.read(32)))        sections: list[TraceSection] = []        for _ in range(header.numProbes):            sections.append(TraceSection(*struct.unpack("IIQ", f.read(16))))        gridSize = header.gridDimX * header.gridDimY * header.gridDimZ        blockSize = header.blockDimX * header.blockDimY * header.blockDimZ        warp_num = blockSize*gridSize // 32        records: dict[str, list[list[NamedTuple]]] = dict()        records["warp_duration"] = []        warp_duration_section = sections[0]        f.seek(warp_duration_section.offset)        per_warprecord_size = warp_duration_section.size        total_buffer_size = warp_num * per_warprecord_size        num_entries = warp_num        print(f"正在读取records, 每个记录大小={per_warprecord_size}, 总记录数={num_entries}")        raw_data = f.read(total_buffer_size)        for i in range(num_entries):            offset = i * per_warprecord_size            if offset + per_warprecord_size > len(raw_data):                print(f"警告: 文件数据不完整，提前结束于第 {i} 条记录。")                break            start, elapsed, warpid = struct.unpack_from("QQQ", raw_data, offset)            records["warp_duration"].append(warp_duration(start=start, elapsed=elapsed, warpid=warpid))        from collections import defaultdict        grouped = defaultdict(list)        for record in records["warp_duration"]:            if record.start > 0or record.elapsed > 0:                grouped[record.warpid].append(record)        # 异常筛选        zero_id_records = [r for r in records["warp_duration"] if r.start == 0and r.elapsed == 0]
        sorted_warpids = sorted(grouped.keys())        print("\n--- Warp Execution Duration Summary ---")        print(f"{'Warp ID':<10}{'Elapsed Cycles (avg)':<20}{'Count':<8}{'Start (first)'}")        for wid in sorted_warpids:            entries = grouped[wid]            avg_elapsed = sum(e.elapsed for e in entries) / len(entries)            valid_starts = [e.start for e in entries]            first_start = min(valid_starts) if valid_starts else0            print(f"{wid:<10}{avg_elapsed:<20.2f}{len(entries):<8}{first_start}")
        if zero_id_records:            print(f"\n注意: 发现了 {len(zero_id_records)} 条无效记录 (start, elapsed均为0)。")            print("可能这些Warps的计时数据没有被成功写回。")
        return header, sections, records
import sysimport numpy as npfrom typing import NamedTuple, List
header, sections, records_map = parse(sys.argv[1]) # filled by path to tracerecords = records_map["warp_duration"]#print(f"check point 1: original data: {records}")

Warp ID    Elapsed Cycles(avg) Count    Start(first)261        138413.00            1        107148575262        138400.00            1        107148577263        138408.00            1        107148579260        138388.00            1        1071485881125       137505.00            1        1071485901126       137552.00            1        1071485921127       137580.00            1        1071485941988       138808.00            1        1071485941989       138846.00            1        1071485961990       138824.00            1        1071485981991       138832.00            1        1071486001124       137521.00            1        107148603..........8629       196405.00            1        1074510097929       195593.00            1        1074510108212       194960.00            1        1074510118352       195822.00            1        1074510127829       195584.00            1        1074510137925       195658.00            1        1074510148208       195607.00            1        1074510158215       195331.00            1        1074510228533       195631.00            1        1074510228632       196351.00            1        107451033.....