如何确定哪些 CUDA 行使用的寄存器最多？

Posted 2023-02-16

【中文标题】如何确定哪些 CUDA 行使用的寄存器最多？

【英文标题】：How to determine which lines of CUDA use the most registers?

【发布时间】：2011-08-19 19:38:16

【问题描述】：

我有一个有点复杂的内核，具有以下统计信息：

ptxas info    : Compiling entry function 'my_kernel' for 'sm_21'
ptxas info    : Function properties for my_kernel
    32 bytes stack frame, 64 bytes spill stores, 40 bytes spill loads
ptxas info    : Used 62 registers, 120 bytes cmem[0], 128 bytes cmem[2], 8 bytes cmem[14], 4 bytes cmem[16]

就寄存器使用而言，我不清楚内核的哪一部分是“高水位线”。内核的性质是，为常量值存根不同部分会导致优化器对后面的部分进行常量折叠，等等（至少看起来是这样，因为我这样做时得到的数字并不多感觉）。

CUDA 分析器对 AFAICT 同样没有帮助，只是告诉我我有注册压力。

有没有办法获得有关寄存器使用的更多信息？我更喜欢某种工具，但我也有兴趣了解直接挖掘已编译的二进制文件，如果需要的话。

编辑：我当然可以采用这种自下而上的方式（即进行实验性代码更改，检查对寄存器使用的影响等），但我宁愿自上而下开始，或者至少得到一些指导从哪里开始自下而上的调查。

【问题讨论】：

您是否在使用 nvcc 编译时专门禁用了设备优化？默认情况下，设备优化已打开。你必须从字面上告诉 nvcc -O0。在我看来，如果你特别禁用优化，编译器根本不应该进行任何代码转换......所以你应该通过存根得到你期望的东西。尽管这可能会为您提供与优化版本不同的 #s 寄存器，但“高水位线”的位置可能在正确的位置......

不，我将继续优化；使用中的寄存器数量只有在优化后才对我感兴趣（如果我不启用优化，编译器会为我完成，我不想花很多精力手动复制寄存器减少）。跨度>

【参考方案1】：

您可以通过像这样编译到带注释的 PTX 来感受编译器输出的复杂性：

nvcc -ptx -Xopencc="-LIST:source=on" branching.cu

它将发出一个 PTX 汇编程序文件，其中包含作为 cmets 的原始 C 代码：

        .entry _Z11branchTest0PfS_S_ (
                .param .u64 __cudaparm__Z11branchTest0PfS_S__a,
                .param .u64 __cudaparm__Z11branchTest0PfS_S__b,
                .param .u64 __cudaparm__Z11branchTest0PfS_S__d)
        
        .reg .u16 %rh<4>;
        .reg .u32 %r<5>;
        .reg .u64 %rd<10>;
        .reg .f32 %f<5>;
        .loc    16      1       0
 //   1  __global__ void branchTest0(float *a, float *b, float *d)
$LDWbegin__Z11branchTest0PfS_S_:
        .loc    16      7       0
 //   3         unsigned int tidx = threadIdx.x + blockDim.x*blockIdx.x;
 //   4         float aval = a[tidx], bval = b[tidx];
 //   5         float z0 = (aval > bval) ? aval : bval;
 //   6  
 //   7         d[tidx] = z0;
        mov.u16         %rh1, %ctaid.x;
        mov.u16         %rh2, %ntid.x;
        mul.wide.u16    %r1, %rh1, %rh2;
        cvt.u32.u16     %r2, %tid.x;
        add.u32         %r3, %r2, %r1;
        cvt.u64.u32     %rd1, %r3;
        mul.wide.u32    %rd2, %r3, 4;
        ld.param.u64    %rd3, [__cudaparm__Z11branchTest0PfS_S__a];
        add.u64         %rd4, %rd3, %rd2;
        ld.global.f32   %f1, [%rd4+0];
        ld.param.u64    %rd5, [__cudaparm__Z11branchTest0PfS_S__b];
        add.u64         %rd6, %rd5, %rd2;
        ld.global.f32   %f2, [%rd6+0];
        max.f32         %f3, %f1, %f2;
        ld.param.u64    %rd7, [__cudaparm__Z11branchTest0PfS_S__d];
        add.u64         %rd8, %rd7, %rd2;
        st.global.f32   [%rd8+0], %f3;
        .loc    16      8       0
 //   8  
        exit;
$LDWend__Z11branchTest0PfS_S_:
         // _Z11branchTest0PfS_S_

请注意，这并没有直接告诉您有关寄存器使用的任何信息，因为 PTX 使用静态单一分配，但它向您展示了汇编程序作为输入给出的内容以及它与原始代码的关系。使用 CUDA 4.0 工具包，您可以将 C 编译为 Fermi 架构的 cubin 文件：

$ nvcc -cubin -arch=sm_20 -Xptxas="-v" branching.cu
ptxas info    : Compiling entry function '_Z11branchTest1PfS_S_' for 'sm_20'
ptxas info    : Function properties for _Z11branchTest1PfS_S_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads

并使用cuobjdump 实用程序反汇编汇编器生成的机器代码。

$ cuobjdump -sass branching.cubin 

code for sm_20
    Function : _Z11branchTest0PfS_S_
/*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
/*0008*/     /*0x94001c042c000000*/     S2R R0, SR_CTAid_X;
/*0010*/     /*0x84009c042c000000*/     S2R R2, SR_Tid_X;
/*0018*/     /*0x10015de218000000*/     MOV32I R5, 0x4;
/*0020*/     /*0x2000dc0320044000*/     IMAD.U32.U32 R3, R0, c [0x0] [0x8], R2;
/*0028*/     /*0x10311c435000c000*/     IMUL.U32.U32.HI R4, R3, 0x4;
/*0030*/     /*0x80319c03200b8000*/     IMAD.U32.U32 R6.CC, R3, R5, c [0x0] [0x20];
/*0038*/     /*0x9041dc4348004000*/     IADD.X R7, R4, c [0x0] [0x24];
/*0040*/     /*0xa0321c03200b8000*/     IMAD.U32.U32 R8.CC, R3, R5, c [0x0] [0x28];
/*0048*/     /*0x00609c8584000000*/     LD.E R2, [R6];
/*0050*/     /*0xb0425c4348004000*/     IADD.X R9, R4, c [0x0] [0x2c];
/*0058*/     /*0xc0329c03200b8000*/     IMAD.U32.U32 R10.CC, R3, R5, c [0x0] [0x30];
/*0060*/     /*0x00801c8584000000*/     LD.E R0, [R8];
/*0068*/     /*0xd042dc4348004000*/     IADD.X R11, R4, c [0x0] [0x34];
/*0070*/     /*0x00201c00081e0000*/     FMNMX R0, R2, R0, !pt;
/*0078*/     /*0x00a01c8594000000*/     ST.E [R10], R0;
/*0080*/     /*0x00001de780000000*/     EXIT;
    ......................................

通常可以从汇编程序回溯到 PTX 并至少大致了解“贪婪”代码部分在哪里。说了这么多，管理寄存器压力是目前 CUDA 编程中比较困难的方面之一。如果/当 NVIDIA 记录其设备代码的 ELF 格式时，我认为合适的代码分析工具对某人来说将是一个很棒的项目。

【讨论】：

你是如何找到 nvcc -ptx -Xopencc="-LIST:source=on" 命令行的？ :D

-Xopencc 不适用于基于 llvm 的 CUDA 版本

@marina.k：是的，但是这个答案是在大约 1 年前写的，当时 CUDA 中还没有基于 LLVM 的工具链。