c - 如何使用 cuobjdump 将 C 代码与 SASS 混合？

Question

我不确定我做错了什么，本质上我想获得可读的程序集，与 C 调用混合。

下面是一些示例代码:

example.cu:

#include <stdio.h>

__global__ void kernel()
{
    unsigned long a, b, c;

    a = 255;
    b = 10;

    c = a + b;
}

int main(void)
{
    cudaFree(0);

    kernel<<<1,1>>>();
    cudaDeviceSynchronize();

    return 0;
}

当我查看 cuobjdump -h(强调我的)时:

--dump-sass (-sass)
Dump assembly for all listed device functions. Cuda source is intermixed with the listed assembly in case option -G was specified to nvcc during compilation, and if the source files can still be found.

我编译它(因此创建了 example.cubin 文件):

nvcc -G -cubin -arch=sm_30 --ptxas-options=-v example.cu

然后我运行:

cuobjdump -sass --function _Z6kernelv example.cubin

输出包含汇编指令，但我在任何地方都看不到 C 代码:

code for sm_30
    Function : _Z6kernelv
.headerflags    @"EF_CUDA_SM30 EF_CUDA_PTX_SM(EF_CUDA_SM30)"
    /*0000*/                   MOV R1, c[0x0][0x44];             /* 0x2800400110005de4 */
    /*0008*/                   ISUB R1, R1, 0x8;                 /* 0x4800c00020105d03 */
    /*0010*/                   S2R R0, SR_LMEMHIOFF;             /* 0x2c000000dc001c04 */
    /*0018*/                   ISETP.GE.AND P0, PT, R1, R0, PT;  /* 0x1b0e00000011dc23 */
    /*0020*/               @P0 BRA 0x30;                         /* 0x40000000200001e7 */
    /*0028*/                   BPT.TRAP 0x1;                     /* 0xd00000000400c007 */
    /*0030*/                   IADD R0, R1, RZ;                  /* 0x48000000fc101c03 */
    /*0038*/                   MOV R2, R0;                       /* 0x2800000000009de4 */
    /*0040*/                   MOV R3, RZ;                       /* 0x28000000fc00dde4 */
    /*0048*/                   MOV R2, R2;                       /* 0x2800000008009de4 */
    /*0050*/                   MOV R3, R3;                       /* 0x280000000c00dde4 */
    /*0058*/                   MOV R4, c[0x0][0x24];             /* 0x2800400090011de4 */
    /*0060*/                   MOV R5, RZ;                       /* 0x28000000fc015de4 */
    /*0068*/                   IADD R2.CC, R2, R4;               /* 0x4801000010209c03 */
    /*0070*/                   IADD.X R3, R3, R5;                /* 0x480000001430dc43 */
    /*0078*/                   MOV32I R4, 0xff;                  /* 0x18000003fc011de2 */
    /*0080*/                   MOV R5, RZ;                       /* 0x28000000fc015de4 */
    /*0088*/                   MOV R4, R4;                       /* 0x2800000010011de4 */

我还没有找到任何选项来明确地告诉 example.cu 所在的位置(虽然它在同一目录中)。具有相同代码的 OTOH Nsight Eclipse Edition 显然能够使用 C 代码显示 SASS(在 Dissassembly 窗口的调试 session 中):

Answer 1

目前，使用 cuobjdump 是不可能的。引用的 cuobjdump 文档/命令行帮助有误。