A CUDA binary (also referred to as cubin) file is an ELF-formatted file which consists of CUDA executable code sections as well as other sections containing symbols, relocators, debug info, etc. By default, the CUDA compiler driver nvcc embeds cubin files into the host executable file. But they can also be generated separately by using the "-cubin" option of nvcc. cubin files are loaded at run time by the CUDA driver API.

CUDA provides two binary utilities for examining and disassembling cubin files and host executables: cuobjdump and nvdisasm. Basically, cuobjdump accepts both cubin files and host binaries while nvdisasm only accepts cubin files; but nvdisasm provides richer output options.

Here's a quick comparison of the two tools:

cuobjdump accepts a single input file each time it's run. The basic usage is as following:

cuobjdump [options] <file>

To disassemble a standalone cubin or cubins embedded in a host executable and show CUDA assembly of the kernels, use the following command:

cuobjdump -sass <input file>

To dump cuda elf sections in human readable format from a cubin file, use the following command:

cuobjdump -elf <cubin file>

To extract ptx text from a host binary, use the following command:

cuobjdump -ptx <host binary>

Here's a sample output of cuobjdump:

$ cuobjdump a.out -ptx -sass
Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
identifier = add.cu

        code for sm_20
                Function : _Z3addPiS_S_
        .headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
        /*0000*/         MOV R1, c[0x1][0x100];  /* 0x2800440400005de4 */
        /*0008*/         MOV R6, c[0x0][0x20];   /* 0x2800400080019de4 */
        /*0010*/         MOV R7, c[0x0][0x24];   /* 0x280040009001dde4 */
        /*0018*/         MOV R2, c[0x0][0x28];   /* 0x28004000a0009de4 */
        /*0020*/         MOV R3, c[0x0][0x2c];   /* 0x28004000b000dde4 */
        /*0028*/         LDU.E R0, [R6];         /* 0x8c00000000601c85 */
        /*0030*/         MOV R4, c[0x0][0x30];   /* 0x28004000c0011de4 */
        /*0038*/         LDU.E R2, [R2];         /* 0x8c00000000209c85 */
        /*0040*/         MOV R5, c[0x0][0x34];   /* 0x28004000d0015de4 */
        /*0048*/         IADD R0, R2, R0;        /* 0x4800000000201c03 */
        /*0050*/         ST.E [R4], R0;          /* 0x9400000000401c85 */
        /*0058*/         EXIT;                   /* 0x8000000000001de7 */
                .............................

Fatbin ptx code:
================
arch = sm_20
code version = [4,0]
producer = cuda
host = linux
compile_size = 64bit
compressed
identifier = add.cu

.version 4.0
.target sm_20
.address_size 64

.visible .entry _Z3addPiS_S_(
.param .u64 _Z3addPiS_S__param_0,
.param .u64 _Z3addPiS_S__param_1,
.param .u64 _Z3addPiS_S__param_2
)
{
.reg .s32 %r<4>;
.reg .s64 %rd<7>;

ld.param.u64 %rd1, [_Z3addPiS_S__param_0];
ld.param.u64 %rd2, [_Z3addPiS_S__param_1];
ld.param.u64 %rd3, [_Z3addPiS_S__param_2];
cvta.to.global.u64 %rd4, %rd3;
cvta.to.global.u64 %rd5, %rd2;
cvta.to.global.u64 %rd6, %rd1;
ldu.global.u32 %r1, [%rd6];
ldu.global.u32 %r2, [%rd5];
add.s32 %r3, %r2, %r1;
st.global.u32 [%rd4], %r3;
ret;
}
        

As shown in the output, the a.out host binary contains cubin and ptx code for sm_20.

To list cubin files in the host binary use -lelf option:

$ cuobjdump a.out -lelf
ELF file    1: add_new.sm_20.cubin
ELF file    2: add_new.sm_30.cubin
ELF file    3: add_old.sm_20.cubin
ELF file    4: add_old.sm_30.cubin
        

To extract all the cubins as files from the host binary use -xelf all option:

$ cuobjdump a.out -xelf all
Extracting ELF file    1: add_new.sm_20.cubin
Extracting ELF file    2: add_new.sm_30.cubin
Extracting ELF file    3: add_old.sm_20.cubin
Extracting ELF file    4: add_old.sm_30.cubin
        

To extract the cubin named add_new.sm_30.cubin:

$ cuobjdump a.out -xelf add_new.sm_30.cubin
Extracting ELF file    1: add_old.sm_20.cubin
        

To extract only the cubins containing _old in their names:

$ cuobjdump a.out -xelf _old
Extracting ELF file    1: add_old.sm_20.cubin
Extracting ELF file    2: add_old.sm_30.cubin
        

You can pass any substring to -xelf and -xptx options. Only the files having the substring in the name will be extracted from the input binary.

To dump common and per function resource usage information:

$ cuobjdump test.cubin -res-usage

Resource usage:
 Common:
  GLOBAL:110 CONSTANT[2]:296 CONSTANT[14]:16
 Function calculate:
  REG:7 STACK:400 SHARED:0 LOCAL:0 CONSTANT[0]:328 TEXTURE:0 SURFACE:0 SAMPLER:0
 Function mysurf_func:
  REG:18 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:444 TEXTURE:0 SURFACE:1 SAMPLER:0
 Function mytexsampler_func:
  REG:42 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:472 TEXTURE:4 SURFACE:0 SAMPLER:1
 Function mysharedfunc:
  REG:30 STACK:0 SHARED:20 LOCAL:0 CONSTANT[0]:192 CONSTANT[16]:440 TEXTURE:0 SURFACE:0 SAMPLER:0
        

Note that value for REG, TEXTURE, SURFACE and SAMPLER denotes the count and for other resources it denotes no. of byte(s) used.

Table 2 contains supported command-line options of cuobjdump, along with a description of what each option does. Each option has a long name and a short name, which can be used interchangeably.

nvdisasm accepts a single input file each time it's run. The basic usage is as following:

nvdisasm [options] <input cubin file>

To get the control flow graph of a kernel, use the following:

nvdisasm -cfg <input cubin file>

Here's a sample output of nvdisasm:

        .headerflags    @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM30 EF_CUDA_PTX_SM(EF_CUDA_SM30) "


//--------------------- .nv.info                  --------------------------
        .section        .nv.info,"",@"SHT_CUDA_INFO "
        .align  4

......

//--------------------- .text._Z4addXPii          --------------------------
        .section        .text._Z4addXPii,"ax",@progbits
        .sectioninfo    @"SHI_REGISTERS=11 "
        .align  4
        .global         _Z4addXPii
        .type           _Z4addXPii,@function
        .size           _Z4addXPii,(.L_19 - _Z4addXPii)
        .other          _Z4addXPii,@"STO_CUDA_ENTRY STV_DEFAULT "
_Z4addXPii:
.text._Z4addXPii:
        /*0008*/                MOV R1, c[0x0][0x44];
        /*0010*/                ISUB R1, R1, 0x8;
        /*0018*/                MOV R0, c[0x0][0x148];
        /*0020*/                IADD R6.CC, R1, c[0x0][0x24];
        /*0028*/                ISETP.LT.AND P0, PT, R0, 0x1, PT;
        /*0030*/                MOV R8, c[0x0][0x140];
        /*0038*/                MOV R9, c[0x0][0x144];
        /*0048*/                IADD.X R7, RZ, RZ;
        /*0050*/                ISUB R10, R6, c[0x0][0x24];
        /*0058*/            @P0 BRA `(.L_2);
        /*0060*/                LD.E R0, [R8];
        /*0068*/                MOV R2, RZ;
        /*0070*/                NOP;
        /*0078*/                NOP;
.L_3:
        /*0088*/                IADD R2, R2, 0x1;
        /*0090*/                MOV R3, R0;
        /*0098*/                IADD R0, R0, 0x1;
        /*00a0*/                ISETP.LT.AND P0, PT, R2, c[0x0][0x148], PT;
        /*00a8*/            @P0 BRA `(.L_3);
        /*00b0*/                IADD R0, R3, 0x1;
        /*00b8*/                ST.E [R8], R0;
.L_2:
        /*00c8*/                S2R R0, SR_TID.X;
        /*00d0*/                ISETP.NE.AND P0, PT, R0, RZ, PT;
        /*00d8*/            @P0 EXIT ;
        /*00e0*/                LD.E R0, [R8];
        /*00e8*/                MOV R4, c[0x0][0xf0];
        /*00f0*/                MOV R5, c[0x0][0xf4];
        /*00f8*/                STL [R10], R0;
        /*0108*/                JCAL `(vprintf);
        /*0110*/                EXIT ;
.L_4:
        /*0118*/                BRA `(.L_4);
.L_19:


//--------------------- SYMBOLS --------------------------

        .type           vprintf,@function
        

nvdisasm is capable of generating control flow of CUDA assembly in the format of DOT graph description language. The output of the control flow from nvdisasm can be directly imported to a DOT graph visualization tool such as Graphviz.

Here's how you can generate a PNG image (cfg.png) of the control flow of the above cubin (a.cubin) with nvdisasm and Graphviz:

nvdisasm -cfg a.cubin | dot -ocfg.png -Tpng

Here's the generated graph:

Figure 1. Control Flow Graph

nvdisasm is capable of showing the register (CC, general and predicate) liveness range information. For each line of CUDA assembly, nvdisasm displays whether a given device register was assigned, accessed, live or re-assigned. It also shows the total number of registers used. This is useful if the user is interested in the life range of any particular register, or register usage in general.

Here's a sample output (left columns are omitted):

                                                          // +------+---------------+-----+
                                                          // | CC   |     GPR       |PRED |
                                                          // |      |    0000000000 |     |
                                                          // | # 01 |  # 0123456789 | # 0 |
                                                          // +------+---------------+-----+
_main10acosParams                                         // |      |               |     |
_main10acosParams,@function                               // |      |               |     |
_main10acosParams,(.L_17 - _Z9acos_main10acosParams)      // |      |               |     |
_main10acosParams,@"STO_CUDA_ENTRY STV_DEFAULT"           // |      |               |     |
                                                          // |      |               |     |
                                                          // |      |               |     |
    MOV R1, c[0x0][0x44];                                 // |      |  1  ^         |     |
    S2R R0, SR_CTAID.X;                                   // |      |  2 ^:         |     |
    S2R R3, SR_TID.X;                                     // |      |  3 :: ^       |     |
    IMAD R3, R0, c[0x0][0x28], R3;                        // |      |  3 v: x       |     |
    MOV R0, c[0x0][0x28];                                 // |      |  3 ^: :       |     |
    ISETP.GE.AND P0, PT, R3, c[0x0][0x150], PT;           // |      |  3 :: v       | 1 ^ |
    IMUL R0, R0, c[0x0][0x34];                            // |      |  3 x: :       | 1 : |
@P0 EXIT;                                                 // |      |  3 :: :       | 1 v |
    MOV32I R8, 0x4;                                       // |      |  4 :: :    ^  |     |
    MOV32I R9, 0x3c94d2e9;                                // |      |  5 :: :    :^ |     |
    NOP;                                                  // |      |  5 :: :    :: |     |
    NOP;                                                  // |      |  5 :: :    :: |     |
    NOP;                                                  // |      |  5 :: :    :: |     |
    NOP;                                                  // |      |  5 :: :    :: |     |
                                                          // |      |  5 :: :    :: |     |
    IMAD R6.CC, R3, R8, c[0x0][0x140];                    // | 1  ^ |  6 :: v  ^ v: |     |
    IMAD.HI.X R7, R3, R8, c[0x0][0x144];                  // | 1  v |  7 :: v  :^v: |     |
    LD.E R2, [R6];                                        // |      |  8 ::^:  vv:: |     |
    FADD.FTZ R4, -|R2|, 1;                                // |      |  7 ::v:^   :: |     |
    FSETP.GT.FTZ.AND P0, PT, |R2|, c[0x2][0x0], PT;       // |      |  7 ::v::   :: | 1 ^ |
    FMUL.FTZ R4, R4, 0.5;                                 // |      |  7 ::::x   :: | 1 : |
    F2F.FTZ.F32.F32 R5, |R2|;                             // |      |  8 ::v::^  :: | 1 : |
    MUFU.RSQ R4, R4;                                      // |      |  8 ::::x:  :: | 1 : |
@P0 MUFU.RCP R5, R4;                                      // |      |  8 ::::v^  :: | 1 v |
    FMUL.FTZ R4, R5, R5;                                  // |      |  8 ::::^v  :: | 1 : |
    IMAD R6.CC, R3, R8, c[0x0][0x148];                    // | 1  ^ |  9 :::v::^ v: | 1 : |
    FFMA.FTZ R7, R4, c[0x2][0x4], R9;                     // | 1  : | 10 ::::v::^:v | 1 : |
    FFMA.FTZ R7, R7, R4, c[0x2][0x8];                     // | 1  : | 10 ::::v::x:: | 1 : |
    FFMA.FTZ R7, R7, R4, c[0x2][0xc];                     // | 1  : | 10 ::::v::x:: | 1 : |
    FFMA.FTZ R7, R7, R4, c[0x2][0x10];                    // | 1  : | 10 ::::v::x:: | 1 : |
    FMUL.FTZ R4, R7, R4;                                  // | 1  : | 10 ::::x::v:: | 1 : |
    IMAD.HI.X R7, R3, R8, c[0x0][0x14c];                  // | 1  v | 10 :::v:::^v: | 1 : |
    FFMA.FTZ R4, R4, R5, R5;                              // |      | 10 ::::xv:::: | 1 : |
    IADD R3, R3, R0;                                      // |      |  9 v::x: :::: | 1 : |
    FADD32I.FTZ R5, -R4, 1.5707963705062866211;           // |      | 10 ::::v^:::: | 1 : |
@P0 FADD.FTZ R5, R4, R4;                                  // |      | 10 ::::v^:::: | 1 v |
    ISETP.LT.AND P0, PT, R3, c[0x0][0x150], PT;           // |      |  9 :::v ::::: | 1 ^ |
    FADD32I.FTZ R4, -R5, 3.1415927410125732422;           // |      | 10 ::::^v:::: | 1 : |
    FCMP.LT.FTZ R2, R4, R5, R2;                           // |      | 10 ::x:vv:::: | 1 : |
    ST.E [R6], R2;                                        // |      |  8 ::v:  vv:: | 1 : |
@P0 BRA `(.L_1);                                          // |      |  5 :: :    :: | 1 v |
    MOV RZ, RZ;                                           // |      |  1  :         |     |
    EXIT;                                                 // |      |  1  :         |     |
                                                          // +......+...............+.....+
    BRA `(.L_2);                                          // |      |               |     |
                                                          // +------+---------------+-----+
                                                          // Legend:
                                                          //     ^       : Register assignment
                                                          //     v       : Register usage
                                                          //     x       : Register usage and reassignment
                                                          //     :       : Register in use
                                                          //     <space> : Register not in use
                                                          //     #       : Number of occupied registers

Table 3 contains the supported command-line options of nvdisasm, along with a description of what each option does. Each option has a long name and a short name, which can be used interchangeably.

(instruction) (destination) (source1), (source2) ...

Table 4 lists valid instructions for the Fermi GPUs.

(instruction) (destination) (source1), (source2) ...

Table 5 lists valid instructions for the Kepler GPUs.

(instruction) (destination) (source1), (source2) ...

Table 6 lists valid instructions for the Maxwell and Pascal GPUs.

(instruction) (destination) (source1), (source2) ...

Table 7 lists valid instructions for the Volta GPUs.

nvprune accepts a single input file each time it's run, emitting a new output file. The basic usage is as following:

nvprune [options] -o <outfile> <infile>

The input file must be either a relocatable host object or static library (not a host executable), and the output file will be the same format.

Either the --arch or --generate-code option must be used to specify the target(s) to keep. All other device code is discarded from the file. The targets can be either a sm_NN arch (cubin) or compute_NN arch (ptx).

For example, the following will prune libcublas_static.a to only contain sm_35 cubin rather than all the targets which normally exist:

nvprune -arch sm_35 libcublas_static.a -o libcublas_static35.a

Note that this means that libcublas_static35.a will not run on any other architecture, so should only be used when you are building for a single architecture.

Table 8 contains supported command-line options of nvprune, along with a description of what each option does. Each option has a long name and a short name, which can be used interchangeably.