cuobjdump与Nsight Compute的SASS差异及地址疑问

问题背景 #

编写的CUDA内核代码：

__global__ void hello_cuda() {
    int a = 10;
    printf("hello from GPU\n");
}

使用Nsight Compute（NCU）查看的SASS内容：

#   Address                   Source
1   00007fe8 82f9ca00        IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] 
2   00007fe8 82f9ca10        UMOV UR4, 0x0 
3   00007fe8 82f9ca20        CS2R R6, SRZ 
4   00007fe8 82f9ca30        UMOV UR5, 0x0 
5   00007fe8 82f9ca40        MOV R4, UR4 
6   00007fe8 82f9ca50        IMAD.U32 R5, RZ, RZ, UR5 
7   00007fe8 82f9ca60        MOV R20, 0x0 
8   00007fe8 82f9ca70        MOV R21, 0x0 
9   00007fe8 82f9ca80        CALL.ABS.NOINC 0x0 
10  00007fe8 82f9ca90        EXIT 
11  00007fe8 82f9caa0        BRA 0x7fe882f9caa0
12  00007fe8 82f9cab0        NOP
13  00007fe8 82f9cac0        NOP
14  00007fe8 82f9cad0        NOP
15  00007fe8 82f9cae0        NOP
16  00007fe8 82f9caf0        NOP

执行cuobjdump -sass saas_source_mapping_trial得到的SASS内容：

code for sm_52
        Function : _Z10hello_cudav
    .headerflags    @"EF_CUDA_SM52 EF_CUDA_PTX_SM(EF_CUDA_SM52)"
                                                           /* 0x001fc400fe2007f6 */
        /*0008*/                   MOV R1, c[0x0][0x20] ;  /* 0x4c98078000870001 */
        /*0010*/                   MOV32I R4, 0x0 ;        /* 0x010000000007f004 */
        /*0018*/                   MOV32I R5, 0x0 ;        /* 0x010000000007f005 */
                                                           /* 0x001ff400fec007f1 */
        /*0028*/                   MOV R6, RZ ;            /* 0x5c9807800ff70006 */
        /*0030*/                   MOV R7, RZ ;            /* 0x5c9807800ff70007 */
        /*0038*/                   JCAL 0x0 ;              /* 0xe220000000000040 */
                                                           /* 0x001ffc00fda007ef */
        /*0048*/                   NOP ;                   /* 0x50b0000000070f00 */
        /*0050*/                   NOP ;                   /* 0x50b0000000070f00 */
        /*0058*/                   EXIT ;                  /* 0xe30000000007000f */
                                                           /* 0x001f8000fc0007ff */
        /*0068*/                   BRA 0x60 ;              /* 0xe2400fffff07000f */
        /*0070*/                   NOP;                    /* 0x50b0000000070f00 */
        /*0078*/                   NOP;                    /* 0x50b0000000070f00 */

疑问解答 #

1. 为何cuobjdump与NCU生成的SASS存在差异？ #

• 工具视角不同：cuobjdump提取的是离线二进制文件中的静态SASS，是编译器输出后存储在可执行文件里的原始内核代码；NCU捕获的是GPU运行时实际执行的SASS，包含CUDA驱动为支持printf这类运行时函数注入的额外包装代码（比如寄存器初始化、运行时服务调用的准备指令）。
• 指令表述差异：两者对同一条硬件指令的命名规则不同。NCU使用更贴近硬件执行的内部指令名（如IMAD.MOV.U32、UMOV），而cuobjdump采用NVIDIA官方文档定义的标准汇编语法（如MOV32I、JCAL），本质是同一指令的不同写法。
• 运行时代码注入：由于内核调用了printf，NCU会显示驱动为了实现GPU端printf而添加的上下文准备代码，而cuobjdump仅展示内核本身的静态编译代码，printf的调用被简化为JCAL指令，不会显示额外的初始化逻辑。

2. NCU的地址是如何获取的？为何与cuobjdump的地址完全不同？ #

• NCU地址是GPU运行时虚拟地址：CUDA驱动加载内核到GPU内存时，会为其分配动态的虚拟地址，这是内核在GPU上实际执行时的内存位置，每次运行可能都不同。
• cuobjdump地址是文件内偏移量：这个地址是内核代码在主机端可执行文件（或.cuobj目标文件）中的相对偏移，仅用于定位静态代码段，和GPU运行时的内存地址没有关联。
• 地址空间完全独立：两者属于不同的地址空间——NCU的地址是GPU虚拟地址空间中的位置，cuobjdump的地址是主机端文件的偏移空间，因此数值完全无关。

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