CUDA二进制工具

cuobjdump、nvdisasm、cu++filt和nvprune的应用说明文档。

1. 概述

本文介绍了cuobjdump、nvdisasm、cu++filt和nvprune这四个适用于Linux(x86、ARM和P9)、Windows、Mac OS和Android平台的CUDA二进制工具。

1.1. 什么是CUDA二进制文件？

CUDA二进制文件（也称为cubin）是一种ELF格式的文件，包含CUDA可执行代码段以及其他包含符号、重定位信息、调试信息等的段。默认情况下，CUDA编译器驱动程序nvcc会将cubin文件嵌入到主机可执行文件中。但也可以通过使用nvcc的“-cubin”选项单独生成它们。cubin文件在运行时由CUDA驱动程序API加载。

注意

有关cubin文件或CUDA编译轨迹的更多详情，请参阅NVIDIA CUDA Compiler Driver NVCC。

1.2. `cuobjdump`与`nvdisasm`的区别

CUDA提供了两个二进制实用工具用于检查和反汇编cubin文件及主机可执行文件：cuobjdump和nvdisasm。基本上，cuobjdump可以处理cubin文件和主机二进制文件，而nvdisasm仅支持cubin文件；但nvdisasm提供了更丰富的输出选项。

以下是两个工具的快速对比：

表1. `cuobjdump` 与 `nvdisasm` 工具对比
	`cuobjdump`	`nvdisasm`
反汇编cubin	是	是
从以下输入文件中提取ptx以及提取并反汇编cubin：主机二进制文件可执行文件目标文件静态库外部fatbinary文件	是	否
控制流分析和输出	否	是
高级显示选项	否	是

1.3. 命令选项类型与表示法

本文档的这一部分提供了以下工具命令行选项的通用详细信息：

cuobjdump
nvdisasm
nvprune

每个命令行选项都有一个长名称和一个短名称，两者可以互换使用。这两种变体的区别在于选项名称前必须使用的连字符数量，即长名称前必须有两个连字符，短名称前必须有一个连字符。例如，-I是--include-path的短名称。长选项适用于构建脚本，其中选项的大小不如描述性值重要，而短选项适用于交互式使用。

上述工具识别三种类型的命令选项：布尔选项、单值选项和列表选项。

布尔选项不需要参数，它们要么在命令行中指定，要么不指定。单值选项最多只能指定一次，而列表选项可以重复。这些选项类型的示例分别是：

Boolean option : nvdisams --print-raw <file>
Single value   : nvdisasm --binary SM70 <file>
List options   : cuobjdump --function "foo,bar,foobar" <file>

单值选项和列表选项必须带有参数，这些参数必须紧跟在选项名称之后，可以通过一个或多个空格或等号字符分隔。当使用单字符短名称（如-I、-l和-L）时，选项值也可以直接跟在选项本身后面，无需通过空格或等号字符分隔。列表选项的各个值可以通过逗号在单个选项实例中分隔，也可以重复该选项，或这两种情况的任意组合。

因此，对于上述两个可能取值的样本选项，以下表示方式是合法的：

-o file
-o=file
-Idir1,dir2 -I=dir3 -I dir4,dir5

对于接受单个值的选项，如果在命令行中多次指定，将采用最右侧的值作为该选项的最终值。在以下示例中，test.bin二进制文件将被反汇编，并假定架构为SM75。

nvdisasm.exe -b SM70 -b SM75 test.bin
nvdisasm warning : incompatible redefinition for option 'binary', the last value of this option was used

对于接受值列表的选项，如果多次指定，这些值会被追加到列表中。如果指定了重复值，它们将被忽略。在下面的示例中，函数foo和bar被视为选项--function的有效值，而重复值foo会被忽略。

cuobjdump --function "foo" --function "bar" --function "foo" -sass  test.cubin

2. cuobjdump

cuobjdump 可从CUDA二进制文件（包括独立文件及嵌入主机二进制文件中的部分）提取信息，并以人类可读格式呈现。cuobjdump 的输出内容包括：每个内核的CUDA汇编代码、CUDA ELF节头、字符串表、重定位表及其他CUDA特定节区。该工具还能从主机二进制文件中提取嵌入的ptx文本。

有关各GPU架构的CUDA汇编指令集列表，请参阅Instruction Set Reference。

2.1. 使用指南

cuobjdump 每次运行时仅接受一个输入文件。基本用法如下：

cuobjdump [options] <file>

要反汇编独立的cubin文件或嵌入在主机可执行文件中的cubin文件并显示内核的CUDA汇编代码，请使用以下命令：

cuobjdump -sass <input file>

要以人类可读的格式从cubin文件中转储cuda elf节，请使用以下命令：

cuobjdump -elf <cubin file>

要从主机二进制文件中提取ptx文本，请使用以下命令：

cuobjdump -ptx <host binary>

以下是cuobjdump的示例输出：

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

code for sm_70
        Function : _Z3addPiS_S_
.headerflags    @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
/*0000*/      IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;  /* 0x00000a00ff017624 */
                                                       /* 0x000fd000078e00ff */
/*0010*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ ;            /* 0x000000fffffff389 */
                                                       /* 0x000fe200000e00ff */
/*0020*/      IMAD.MOV.U32 R2, RZ, RZ, c[0x0][0x160] ; /* 0x00005800ff027624 */
                                                       /* 0x000fe200078e00ff */
/*0030*/      MOV R3, c[0x0][0x164] ;                  /* 0x0000590000037a02 */
                                                       /* 0x000fe20000000f00 */
/*0040*/      IMAD.MOV.U32 R4, RZ, RZ, c[0x0][0x168] ; /* 0x00005a00ff047624 */
                                                       /* 0x000fe200078e00ff */
/*0050*/      MOV R5, c[0x0][0x16c] ;                  /* 0x00005b0000057a02 */
                                                       /* 0x000fcc0000000f00 */
/*0060*/      LDG.E.SYS R2, [R2] ;                     /* 0x0000000002027381 */
                                                       /* 0x000ea800001ee900 */
/*0070*/      LDG.E.SYS R5, [R4] ;                     /* 0x0000000004057381 */
                                                       /* 0x000ea200001ee900 */
/*0080*/      IMAD.MOV.U32 R6, RZ, RZ, c[0x0][0x170] ; /* 0x00005c00ff067624 */
                                                       /* 0x000fe200078e00ff */
/*0090*/      MOV R7, c[0x0][0x174] ;                  /* 0x00005d0000077a02 */
                                                       /* 0x000fe40000000f00 */
/*00a0*/      IADD3 R9, R2, R5, RZ ;                   /* 0x0000000502097210 */
                                                       /* 0x004fd00007ffe0ff */
/*00b0*/      STG.E.SYS [R6], R9 ;                     /* 0x0000000906007386 */
                                                       /* 0x000fe2000010e900 */
/*00c0*/      EXIT ;                                   /* 0x000000000000794d */
                                                       /* 0x000fea0003800000 */
/*00d0*/      BRA 0xd0;                                /* 0xfffffff000007947 */
                                                       /* 0x000fc0000383ffff */
/*00e0*/      NOP;                                     /* 0x0000000000007918 */
                                                       /* 0x000fc00000000000 */
/*00f0*/      NOP;                                     /* 0x0000000000007918 */
                                                       /* 0x000fc00000000000 */
        .......................

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

.version 7.0
.target sm_70
.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;
ld.global.u32 %r1, [%rd6];
ld.global.u32 %r2, [%rd5];
add.s32 %r3, %r2, %r1;
st.global.u32 [%rd4], %r3;
ret;
}

如输出所示，a.out主机二进制文件包含针对sm_70架构的cubin和ptx代码。

要列出主机二进制文件中的cubin文件，请使用-lelf选项：

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

要从主机二进制文件中提取所有cubin文件，请使用-xelf all选项：

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

要提取名为add_new.sm_70.cubin的cubin文件：

$ cuobjdump a.out -xelf add_new.sm_70.cubin
Extracting ELF file    1: add_new.sm_70.cubin

仅提取名称中包含_old的cubin文件：

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

你可以向-xelf和-xptx选项传递任何子字符串。只有名称中包含该子字符串的文件才会从输入二进制文件中提取出来。

转储通用及每个函数的资源使用信息：

$ cuobjdump test.cubin -res-usage

Resource usage:
 Common:
  GLOBAL:56 CONSTANT[3]:28
 Function calculate:
  REG:24 STACK:8 SHARED:0 LOCAL:0 CONSTANT[0]:472 CONSTANT[2]:24 TEXTURE:0 SURFACE:0 SAMPLER:0
 Function mysurf_func:
  REG:38 STACK:8 SHARED:4 LOCAL:0 CONSTANT[0]:532 TEXTURE:8 SURFACE:7 SAMPLER:0
 Function mytexsampler_func:
  REG:42 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:472 TEXTURE:4 SURFACE:0 SAMPLER:1

请注意，REG、TEXTURE、SURFACE和SAMPLER的值表示数量，而对于其他资源则表示使用的字节数。

2.2. 命令行选项

表2列出了cuobjdump支持的命令行选项，以及每个选项功能的描述。每个选项都有长名称和短名称，可以互换使用。

表2. `cuobjdump` 命令行选项
选项（长格式）	选项（短格式）	描述
`--all-fatbin`	`-all`	转储所有fatbin段。默认情况下，仅转储可执行fatbin的内容（如果存在），如果没有可执行fatbin，则转储可重定位fatbin。
`--dump-elf`	`-elf`	转储ELF对象段。
`--dump-elf-symbols`	`-symbols`	转储ELF符号名称。
`--dump-ptx`	`-ptx`	转储所有列出设备函数的PTX代码。
`--dump-sass`	`-sass`	转储单个cubin文件或嵌入在二进制文件中的所有cubin文件的CUDA汇编代码。
`--dump-resource-usage`	`-res-usage`	转储每个ELF文件的资源使用情况。有助于在一个地方获取所有资源使用信息。
`--extract-elf <部分文件名 >,...`	`-xelf`	提取包含<部分文件名>的ELF文件名称并保存为文件。使用`all`可提取所有文件。要获取ELF文件列表，请使用-lelf选项。适用于主机可执行文件/对象/库以及外部fatbin。所有`dump`和`list`选项在使用此选项时将被忽略。
`--extract-ptx file name>,...`	`-xptx`	提取包含<部分文件名>的PTX文件名称并保存为文件。使用`all`可提取所有文件。要获取PTX文件列表，请使用-lptx选项。适用于主机可执行文件/对象/库和外部fatbin。所有`dump`和`list`选项在此选项下将被忽略。
`--extract-text file name>,...`	`-xtext`	提取包含<部分文件名>的文本二进制编码文件名称并保存为文件。使用'all'可提取所有文件。要获取文本二进制编码列表，请使用-ltext选项。所有'dump'和'list'选项在此情况下将被忽略。
`--function name>,...`	`-fun`	指定必须转储其胖二进制结构的设备函数名称。
`--function-index index>,...`	`-findex`	指定必须转储其fat二进制结构的函数的符号表索引。
`--gpu-architecture architecture name>`	`-arch`	指定需要转储信息的GPU架构。此选项允许的值为：`sm_50`, `sm_52`, `sm_53`, `sm_60`, `sm_61`, `sm_62`, `sm_70`, `sm_72`, `sm_75`, `sm_80`, `sm_86`, `sm_87`, `sm_89`, `sm_90`, `sm_90a`, `sm_100`, `sm_100a`, `sm_120`, `sm_120a`。
`--help`	`-h`	打印此工具的帮助信息。
`--list-elf`	`-lelf`	列出fatbin中所有可用的ELF文件。适用于主机可执行文件/对象/库以及外部fatbin。使用此标志时，所有其他选项将被忽略。之后可通过-xelf选项选择特定的ELF文件。
`--list-ptx`	`-lptx`	列出fatbin中所有可用的PTX文件。适用于主机可执行文件/对象/库以及外部fatbin。使用此标志时，所有其他选项将被忽略。后续可通过-xptx选项选择特定的PTX文件。
`--list-text`	`-ltext`	列出fatbin中所有可用的文本二进制函数名称。使用此标志时，其他所有选项都将被忽略。之后可以通过-xtext选项选择特定函数。
`--options-file ,...`	`-optf`	从指定文件包含命令行选项。
`--sort-functions`	`-sort`	转储sass时对函数进行排序。
`--version`	`-V`	打印此工具的版本信息。

3. nvdisasm

nvdisasm 从独立的cubin文件中提取信息，并以人类可读的格式呈现。nvdisasm的输出包括每个内核的CUDA汇编代码、ELF数据段列表以及其他CUDA特定段。输出样式和选项通过nvdisasm命令行选项控制。nvdisasm还会进行控制流分析，标注跳转/分支目标，使输出更易于阅读。

注意

nvdisasm 需要完整的重定位信息来进行控制流分析。如果CUDA二进制文件中缺少此信息，可以使用nvdisasm选项-ndf关闭控制流分析，或者使用ptxas和nvlink选项-preserve-relocs重新生成cubin文件。

有关各GPU架构的CUDA汇编指令集列表，请参阅Instruction Set Reference。

3.1. 使用指南

nvdisasm 每次运行时接受单个输入文件。基本用法如下：

nvdisasm [options] <input cubin file>

以下是nvdisasm的示例输出：

    .headerflags    @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM70
                      EF_CUDA_VIRTUAL_SM(EF_CUDA_SM70)"
    .elftype        @"ET_EXEC"

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

......

//--------------------- .text._Z9acos_main10acosParams --------------------------
    .section    .text._Z9acos_main10acosParams,"ax",@progbits
    .sectioninfo    @"SHI_REGISTERS=14"
    .align    128
        .global     _Z9acos_main10acosParams
        .type       _Z9acos_main10acosParams,@function
        .size       _Z9acos_main10acosParams,(.L_21 - _Z9acos_main10acosParams)
        .other      _Z9acos_main10acosParams,@"STO_CUDA_ENTRY STV_DEFAULT"
_Z9acos_main10acosParams:
.text._Z9acos_main10acosParams:
        /*0000*/               MOV R1, c[0x0][0x28] ;
        /*0010*/               NOP;
        /*0020*/               S2R R0, SR_CTAID.X ;
        /*0030*/               S2R R3, SR_TID.X ;
        /*0040*/               IMAD R0, R0, c[0x0][0x0], R3 ;
        /*0050*/               ISETP.GE.AND P0, PT, R0, c[0x0][0x170], PT ;
        /*0060*/           @P0 EXIT ;
.L_1:
        /*0070*/               MOV R11, 0x4 ;
        /*0080*/               IMAD.WIDE R2, R0, R11, c[0x0][0x160] ;
        /*0090*/               LDG.E.SYS R2, [R2] ;
        /*00a0*/               MOV R7, 0x3d53f941 ;
        /*00b0*/               FADD.FTZ R4, |R2|.reuse, -RZ ;
        /*00c0*/               FSETP.GT.FTZ.AND P0, PT, |R2|.reuse, 0.5699, PT ;
        /*00d0*/               FSETP.GEU.FTZ.AND P1, PT, R2, RZ, PT ;
        /*00e0*/               FADD.FTZ R5, -R4, 1 ;
        /*00f0*/               IMAD.WIDE R2, R0, R11, c[0x0][0x168] ;
        /*0100*/               FMUL.FTZ R5, R5, 0.5 ;
        /*0110*/           @P0 MUFU.SQRT R4, R5 ;
        /*0120*/               MOV R5, c[0x0][0x0] ;
        /*0130*/               IMAD R0, R5, c[0x0][0xc], R0 ;
        /*0140*/               FMUL.FTZ R6, R4, R4 ;
        /*0150*/               FFMA.FTZ R7, R6, R7, 0.018166976049542427063 ;
        /*0160*/               FFMA.FTZ R7, R6, R7, 0.046756859868764877319 ;
        /*0170*/               FFMA.FTZ R7, R6, R7, 0.074846573173999786377 ;
        /*0180*/               FFMA.FTZ R7, R6, R7, 0.16667014360427856445 ;
        /*0190*/               FMUL.FTZ R7, R6, R7 ;
        /*01a0*/               FFMA.FTZ R7, R4, R7, R4 ;
        /*01b0*/               FADD.FTZ R9, R7, R7 ;
        /*01c0*/          @!P0 FADD.FTZ R9, -R7, 1.5707963705062866211 ;
        /*01d0*/               ISETP.GE.AND P0, PT, R0, c[0x0][0x170], PT ;
        /*01e0*/          @!P1 FADD.FTZ R9, -R9, 3.1415927410125732422 ;
        /*01f0*/               STG.E.SYS [R2], R9 ;
        /*0200*/          @!P0 BRA `(.L_1) ;
        /*0210*/               EXIT ;
.L_2:
        /*0220*/               BRA `(.L_2);
.L_21:

要获取内核的控制流图，请使用以下方法：

nvdisasm -cfg <input cubin file>

nvdisasm 能够以DOT图形描述语言格式生成CUDA汇编的控制流。从nvdisasm输出的控制流可以直接导入到诸如Graphviz等DOT图形可视化工具中。

以下是使用nvdisasm和Graphviz生成上述cubin文件(a.cubin)控制流PNG图像(cfg.png)的方法：

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

这是生成的图表：

要使用nvdisasm和Graphviz生成上述cubin文件(a.cubin)的基本块控制流的PNG图像(bbcfg.png)：

nvdisasm -bbcfg a.cubin | dot -obbcfg.png -Tpng

这是生成的图表：

Basic Block Control Flow Graph — 基础块控制流图

nvdisasm能够显示寄存器（通用和谓词）的生命周期范围信息。对于每一行CUDA汇编代码，nvdisasm会显示给定的设备寄存器是否被分配、访问、存活或重新分配。它还会显示使用的寄存器总数。如果用户对特定寄存器的生命周期范围或一般寄存器使用情况感兴趣，这将非常有用。

以下是一个示例输出（为简洁起见，输出内容已精简）：

                                                      // +-----------------+------+
                                                      // |      GPR        | PRED |
                                                      // |                 |      |
                                                      // |                 |      |
                                                      // |    000000000011 |      |
                                                      // |  # 012345678901 | # 01 |
                                                      // +-----------------+------+
    .global acos                                      // |                 |      |
    .type   acos,@function                            // |                 |      |
    .size   acos,(.L_21 - acos)                       // |                 |      |
    .other  acos,@"STO_CUDA_ENTRY STV_DEFAULT"        // |                 |      |
acos:                                                 // |                 |      |
.text.acos:                                           // |                 |      |
    MOV R1, c[0x0][0x28] ;                            // |  1  ^           |      |
    NOP;                                              // |  1  ^           |      |
    S2R R0, SR_CTAID.X ;                              // |  2 ^:           |      |
    S2R R3, SR_TID.X ;                                // |  3 :: ^         |      |
    IMAD R0, R0, c[0x0][0x0], R3 ;                    // |  3 x: v         |      |
    ISETP.GE.AND P0, PT, R0, c[0x0][0x170], PT ;      // |  2 v:           | 1 ^  |
@P0 EXIT ;                                            // |  2 ::           | 1 v  |
.L_1:                                                 // |  2 ::           |      |
     MOV R11, 0x4 ;                                   // |  3 ::         ^ |      |
     IMAD.WIDE R2, R0, R11, c[0x0][0x160] ;           // |  5 v:^^       v |      |
     LDG.E.SYS R2, [R2] ;                             // |  4 ::^        : |      |
     MOV R7, 0x3d53f941 ;                             // |  5 :::    ^   : |      |
     FADD.FTZ R4, |R2|.reuse, -RZ ;                   // |  6 ::v ^  :   : |      |
     FSETP.GT.FTZ.AND P0, PT, |R2|.reuse, 0.5699, PT; // |  6 ::v :  :   : | 1 ^  |
     FSETP.GEU.FTZ.AND P1, PT, R2, RZ, PT ;           // |  6 ::v :  :   : | 2 :^ |
     FADD.FTZ R5, -R4, 1 ;                            // |  6 ::  v^ :   : | 2 :: |
     IMAD.WIDE R2, R0, R11, c[0x0][0x168] ;           // |  8 v:^^:: :   v | 2 :: |
     FMUL.FTZ R5, R5, 0.5 ;                           // |  5 ::  :x :     | 2 :: |
 @P0 MUFU.SQRT R4, R5 ;                               // |  5 ::  ^v :     | 2 v: |
     MOV R5, c[0x0][0x0] ;                            // |  5 ::  :^ :     | 2 :: |
     IMAD R0, R5, c[0x0][0xc], R0 ;                   // |  5 x:  :v :     | 2 :: |
     FMUL.FTZ R6, R4, R4 ;                            // |  5 ::  v ^:     | 2 :: |
     FFMA.FTZ R7, R6, R7, 0.018166976049542427063 ;   // |  5 ::  : vx     | 2 :: |
     FFMA.FTZ R7, R6, R7, 0.046756859868764877319 ;   // |  5 ::  : vx     | 2 :: |
     FFMA.FTZ R7, R6, R7, 0.074846573173999786377 ;   // |  5 ::  : vx     | 2 :: |
     FFMA.FTZ R7, R6, R7, 0.16667014360427856445 ;    // |  5 ::  : vx     | 2 :: |
     FMUL.FTZ R7, R6, R7 ;                            // |  5 ::  : vx     | 2 :: |
     FFMA.FTZ R7, R4, R7, R4 ;                        // |  4 ::  v  x     | 2 :: |
     FADD.FTZ R9, R7, R7 ;                            // |  4 ::     v ^   | 2 :: |
@!P0 FADD.FTZ R9, -R7, 1.5707963705062866211 ;        // |  4 ::     v ^   | 2 v: |
     ISETP.GE.AND P0, PT, R0, c[0x0][0x170], PT ;     // |  3 v:       :   | 2 ^: |
@!P1 FADD.FTZ R9, -R9, 3.1415927410125732422 ;        // |  3 ::       x   | 2 :v |
     STG.E.SYS [R2], R9 ;                             // |  3 ::       v   | 1 :  |
@!P0 BRA `(.L_1) ;                                    // |  2 ::           | 1 v  |
     EXIT ;                                           // |  1  :           |      |
.L_2:                                                 // +.................+......+
     BRA `(.L_2);                                     // |                 |      |
.L_21:                                                // +-----------------+------+
                                                      // Legend:
                                                      //     ^       : Register assignment
                                                      //     v       : Register usage
                                                      //     x       : Register usage and reassignment
                                                      //     :       : Register in use
                                                      //     <space> : Register not in use
                                                      //     #       : Number of occupied registers

nvdisasm 能够显示CUDA源文件的行号信息，这对调试非常有用。

要获取内核的行信息，请使用以下方法：

nvdisasm -g <input cubin file>

以下是使用nvdisasm -g命令的内核示例输出：

//--------------------- .text._Z6kernali          --------------------------
        .section        .text._Z6kernali,"ax",@progbits
        .sectioninfo    @"SHI_REGISTERS=24"
        .align  128
        .global         _Z6kernali
        .type           _Z6kernali,@function
        .size           _Z6kernali,(.L_4 - _Z6kernali)
        .other          _Z6kernali,@"STO_CUDA_ENTRY STV_DEFAULT"
_Z6kernali:
.text._Z6kernali:
        /*0000*/                   MOV R1, c[0x0][0x28] ;
        /*0010*/                   NOP;
    //## File "/home/user/cuda/sample/sample.cu", line 25
        /*0020*/                   MOV R0, 0x160 ;
        /*0030*/                   LDC R0, c[0x0][R0] ;
        /*0040*/                   MOV R0, R0 ;
        /*0050*/                   MOV R2, R0 ;
    //## File "/home/user/cuda/sample/sample.cu", line 26
        /*0060*/                   MOV R4, R2 ;
        /*0070*/                   MOV R20, 32@lo((_Z6kernali + .L_1@srel)) ;
        /*0080*/                   MOV R21, 32@hi((_Z6kernali + .L_1@srel)) ;
        /*0090*/                   CALL.ABS.NOINC `(_Z3fooi) ;
.L_1:
        /*00a0*/                   MOV R0, R4 ;
        /*00b0*/                   MOV R4, R2 ;
        /*00c0*/                   MOV R2, R0 ;
        /*00d0*/                   MOV R20, 32@lo((_Z6kernali + .L_2@srel)) ;
        /*00e0*/                   MOV R21, 32@hi((_Z6kernali + .L_2@srel)) ;
        /*00f0*/                   CALL.ABS.NOINC `(_Z3bari) ;
.L_2:
        /*0100*/                   MOV R4, R4 ;
        /*0110*/                   IADD3 R4, R2, R4, RZ ;
        /*0120*/                   MOV R2, 32@lo(arr) ;
        /*0130*/                   MOV R3, 32@hi(arr) ;
        /*0140*/                   MOV R2, R2 ;
        /*0150*/                   MOV R3, R3 ;
        /*0160*/                   ST.E.SYS [R2], R4 ;
    //## File "/home/user/cuda/sample/sample.cu", line 27
        /*0170*/                   ERRBAR ;
        /*0180*/                   EXIT ;
.L_3:
        /*0190*/                   BRA `(.L_3);
.L_4:

nvdisasm 能够显示行号信息以及额外的函数内联信息（如果有的话）。在没有函数内联的情况下，其输出与使用 nvdisasm -g 命令时的输出相同。

以下是使用nvdisasm -gi命令的内核示例输出：

//--------------------- .text._Z6kernali          --------------------------
    .section    .text._Z6kernali,"ax",@progbits
    .sectioninfo    @"SHI_REGISTERS=16"
    .align    128
        .global         _Z6kernali
        .type           _Z6kernali,@function
        .size           _Z6kernali,(.L_18 - _Z6kernali)
        .other          _Z6kernali,@"STO_CUDA_ENTRY STV_DEFAULT"
_Z6kernali:
.text._Z6kernali:
        /*0000*/                   IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*0010*/                   UMOV UR4, 32@lo(arr) ;
        /*0020*/                   UMOV UR5, 32@hi(arr) ;
        /*0030*/                   IMAD.U32 R2, RZ, RZ, UR4 ;
        /*0040*/                   MOV R3, UR5 ;
        /*0050*/                   ULDC.64 UR4, c[0x0][0x118] ;
    //## File "/home/user/cuda/inline.cu", line 10 inlined at "/home/user/cuda/inline.cu", line 17
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*0060*/                   LDG.E R4, [R2.64] ;
        /*0070*/                   LDG.E R5, [R2.64+0x4] ;
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*0080*/                   LDG.E R0, [R2.64+0x8] ;
    //## File "/home/user/cuda/inline.cu", line 23
        /*0090*/                   UMOV UR6, 32@lo(ans) ;
        /*00a0*/                   UMOV UR7, 32@hi(ans) ;
    //## File "/home/user/cuda/inline.cu", line 10 inlined at "/home/user/cuda/inline.cu", line 17
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*00b0*/                   IADD3 R7, R4, c[0x0][0x160], RZ ;
    //## File "/home/user/cuda/inline.cu", line 23
        /*00c0*/                   IMAD.U32 R4, RZ, RZ, UR6 ;
    //## File "/home/user/cuda/inline.cu", line 10 inlined at "/home/user/cuda/inline.cu", line 17
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*00d0*/                   IADD3 R9, R5, c[0x0][0x160], RZ ;
    //## File "/home/user/cuda/inline.cu", line 23
        /*00e0*/                   MOV R5, UR7 ;
    //## File "/home/user/cuda/inline.cu", line 10 inlined at "/home/user/cuda/inline.cu", line 17
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*00f0*/                   IADD3 R11, R0.reuse, c[0x0][0x160], RZ ;
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*0100*/                   IMAD.IADD R13, R0, 0x1, R7 ;
    //## File "/home/user/cuda/inline.cu", line 10 inlined at "/home/user/cuda/inline.cu", line 17
    //## File "/home/user/cuda/inline.cu", line 17 inlined at "/home/user/cuda/inline.cu", line 23
    //## File "/home/user/cuda/inline.cu", line 23
        /*0110*/                   STG.E [R2.64+0x4], R9 ;
        /*0120*/                   STG.E [R2.64], R7 ;
        /*0130*/                   STG.E [R2.64+0x8], R11 ;
    //## File "/home/user/cuda/inline.cu", line 23
        /*0140*/                   STG.E [R4.64], R13 ;
    //## File "/home/user/cuda/inline.cu", line 24
        /*0150*/                   EXIT ;
.L_3:
        /*0160*/                   BRA (.L_3);
.L_18:

nvdisasm 可以生成JSON格式的反汇编代码。

有关JSON格式的详细信息，请参阅附录。

要获取JSON格式的反汇编代码，请使用以下方法：

nvdisasm -json <input cubin file>

nvdisasm -json 的输出将是压缩格式。下面的示例是经过美化后的结果：

[
    {
        "ELF": {
            "layout-id": 4,
            "ei_osabi": 51,
            "ei_abiversion": 7
        },
        "SM": {
            "version": {
                "major": 9,
                "minor": 0
            }
        },
        "SchemaVersion": {
            "major": 12,
            "minor": 8,
            "revision": 0
        },
        "Producer": "nvdisasm V12.8.14 Build r570_00.r12.8/compiler.35033008_0",
        "Description": ""
    },
    [
        {
            "function-name": "foo",
            "start": 0,
            "length": 96,
            "other-attributes": [],
            "sass-instructions": [
                {
                    "opcode": "LDC",
                    "operands": "R1,c[0x0][0x28]"
                },
                {
                    "opcode": "MOV",
                    "operands": "R6,0x60"
                },
                {
                    "opcode": "ISETP.NE.U32.AND",
                    "operands": "P0,PT,R1,0x1,PT"
                },
                {
                    "opcode": "CALL.REL.NOINC",
                    "operands": "`(bar)",
                    "other-attributes": {
                        "control-flow": "True"
                    }
                },
                {
                    "opcode": "MOV",
                    "operands": "R8,R7"
                },
                {
                    "opcode": "EXIT",
                    "other-attributes": {
                        "control-flow": "True"
                    }
                }
            ]
        },
        {
            "function-name": "bar",
            "start": 96,
            "length": 32,
            "other-attributes": [],
            "sass-instructions": [
                {
                    "opcode": "STS.128",
                    "operands": "[UR5+0x400],RZ"
                },
                {
                    "opcode": "RET.REL.NODEC",
                    "operands": "R18,`(foo)",
                    "other-attributes": {
                        "control-flow": "True"
                    }
                }
            ]
        }
    ]
]

3.2. 命令行选项

表3列出了nvdisasm支持的命令行选项，以及每个选项功能的描述。每个选项都有长名称和短名称，可以互换使用。

表3. nvdisasm命令行选项
选项（长格式）	选项（短格式）	描述
`--base-address`	`-base`	指定要反汇编的映像的逻辑基地址。此选项仅在反汇编原始指令二进制文件时有效（参见选项`--binary`），在反汇编Elf文件时将被忽略。默认值：0。
`--binary`	`-b`	当指定此选项时，输入文件将被视为包含原始指令二进制文件，即指令内存中出现的一系列二进制指令编码。此选项的值必须是原始二进制文件所声明的架构。此选项允许的值为：`SM50`, `SM52`, `SM53`, `SM60`, `SM61`, `SM62`, `SM70`, `SM72`, `SM75`, `SM80`, `SM86`, `SM87`, `SM89`, `SM90`, `SM90a`, `SM100`, `SM100a`, `SM100`, `SM100a`, `SM120`, `SM120a`。
`--cuda-function-index <符号索引>,...`	`-fun`	将输出限制为具有给定索引符号所表示的CUDA函数。对于给定符号的CUDA函数是其所属的代码段。此限制仅适用于可执行段；所有其他段仍将被输出。
`--emit-json`	`-json`	以JSON格式打印反汇编代码。这可以与选项`--binary` 和`--cuda-function-index index>,...`一起使用。关于JSON格式的详细信息，请参阅附录。但此功能与选项`--print-life-ranges`、`--life-range-mode`、`--output-control-flow-graph`以及`--output-control-flow-graph-with-basic-blocks`不兼容。
`--help`	`-h`	打印此工具的帮助信息。
`--life-range-mode`	`-lrm`	此选项隐含`--print-life-ranges`选项，用于确定如何打印寄存器存活范围信息。`count`：完全不打印，仅保留#列（存活寄存器数量）；`wide`：为可读性而分散列（默认）；`narrow`：每个寄存器使用单字符列，节省表格宽度。此选项允许的值：`count`、`narrow`、`wide`。
`--no-dataflow`	`-ndf`	在反汇编后禁用数据流分析器。通常启用数据流分析是为了执行分支堆栈分析，并通过推断的分支目标标签来标注所有通过GPU分支堆栈跳转的指令。但当输入的nvelf/cubin不满足某些限制条件时，该分析可能偶尔会失败。
`--no-vliw`	`-novliw`	常规模式；以普通语法而非VLIW语法反汇编配对指令。
`--options-file <file>,...`	`-optf`	从指定文件包含命令行选项。
`--output-control-flow-graph`	`-cfg`	当指定输出控制流图时，每个节点代表一个超块，格式兼容graphviz工具（如dot）。
`--output-control-flow-graph-with-basic-blocks`	`-bbcfg`	当指定输出控制流图时，每个节点代表一个基本块，格式兼容graphviz工具（如dot）。
`--打印代码`	`-c`	仅打印代码部分。
`--print-instr-offsets-cfg`	`-poff`	当指定时，在控制流图中打印指令偏移量。此选项应与`--output-control-flow-graph`或`--output-control-flow-graph-with-basic-blocks`一起使用。
`--print-instruction-encoding`	`-hex`	当指定时，在每条反汇编操作后打印编码字节。
`--print-life-ranges`	`-plr`	在生成的汇编代码末尾列中打印寄存器生命周期范围信息。
`--打印行信息`	`-g`	如果存在.debug_line节，则使用从中获取的源代码行信息注释反汇编。
`--print-line-info-inline`	`-gi`	使用从.debug_line部分获取的源代码行信息以及函数内联信息（如果存在）来注释反汇编代码。
`--print-line-info-ptx`	`-gp`	如果存在，使用从.nv_debug_line_sass部分获取的源代码行信息来注释反汇编。
`--print-raw`	`-raw`	打印反汇编代码，不进行任何美化处理。
`--separate-functions`	`-sf`	将与函数符号对应的代码通过一些空行分隔开，使它们在打印的反汇编中更加突出。
`--version`	`-V`	打印此工具的版本信息。

4. 指令集参考

本节包含NVIDIA^® GPU架构的指令集参考。

4.1. Maxwell与Pascal指令集

Maxwell（计算能力5.x）和Pascal（计算能力6.x）架构具有以下指令集格式：

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

有效的目标位置和源位置包括：

寄存器RX
用于特殊系统控制寄存器的SRX
条件寄存器的PX
c[X][Y] 表示常量内存

表4列出了Maxwell和Pascal GPU的有效指令集。

表4. Maxwell与Pascal指令集
操作码	描述
浮点指令
FADD	FP32加法
FCHK	单精度浮点除法范围检查
FCMP	FP32 与零比较并选择源
FFMA	FP32 融合乘加运算
FMNMX	FP32 最小值/最大值
FMUL	FP32乘法
FSET	FP32比较并设置
FSETP	FP32比较并设置谓词
FSWZADD	用于FSWZ仿真的FP32加法运算
MUFU	多功能操作
RRO	范围缩减运算符FP
DADD	双精度浮点数加法
DFMA	FP64 融合乘加运算
DMNMX	FP64 最小值/最大值
DMUL	双精度浮点数乘法
DSET	FP64比较并设置
DSETP	FP64比较并设置谓词
HADD2	FP16加法
HFMA2	FP16融合乘加运算
HMUL2	FP16乘法运算
HSET2	FP16比较并设置
HSETP2	FP16比较并设置谓词
整数指令
BFE	位域提取
BFI	位域插入
FLO	查找前导1
IADD	整数加法
IADD3	三输入整数加法
ICMP	整数与零比较并选择源
IMAD	整数乘加运算
IMADSP	提取的整数乘加运算。
IMNMX	整数最小值/最大值
IMUL	整数乘法
ISCADD	缩放整数加法
ISET	整数比较并设置
ISETP	整数比较与设置谓词
LEA	计算有效地址
LOP	逻辑运算
LOP3	三输入逻辑运算
POPC	人口计数
SHF	漏斗位移
SHL	左移
SHR	右移
XMAD	整数短乘加运算
转换说明
F2F	浮点数到浮点数转换
F2I	浮点数转整数转换
I2F	整数到浮点数转换
I2I	整数到整数转换
移动指令
MOV	移动
PRMT	寄存器对置换
SEL	带谓词的选择源
SHFL	线程束级寄存器洗牌
谓词/CC指令
CSET	测试条件代码并设置
CSETP	测试条件代码并设置谓词
PSET	组合谓词与集合
PSETP	组合谓词并设置谓词
P2R	将谓词寄存器移动到寄存器
R2P	将寄存器移动到谓词/条件码寄存器
纹理指令
TEX	纹理获取
TLD	纹理加载
TLD4	纹理加载4
TXQ	纹理查询
TEXS	使用标量/非vec4源/目标进行纹理提取
TLD4S	使用标量/非vec4源/目标的纹理加载4
TLDS	使用标量/非vec4源/目标的纹理加载
计算加载/存储指令
LD	从通用内存加载
LDC	加载常量
LDG	从全局内存加载
LDL	在本地内存窗口内加载
LDS	共享内存窗口内的本地访问
ST	存储到通用内存
STG	存储到全局内存
STL	存储到本地内存
STS	存储到共享内存
ATOM	通用内存原子操作
ATOMS	共享内存原子操作
RED	通用内存上的归约操作
CCTL	缓存控制
CCTLL	缓存控制
MEMBAR	内存屏障
CCTLT	纹理缓存控制
表面内存指令
SUATOM	表面内存原子操作
SULD	表面负载
SURED	表面内存上的归约操作
SUST	表面存储
控制指令
BRA	相对分支
BRX	相对间接分支
JMP	绝对跳转
JMX	绝对间接跳转
SSY	设置同步点
SYNC	在条件分支后同步线程
CAL	相对调用
JCAL	绝对调用
PRET	子程序预返回
RET	从子程序返回
BRK	中断
PBK	预中断
CONT	继续
PCNT	预继续
EXIT	退出程序
PEXIT	预退出
BPT	断点/陷阱
其他指令
NOP	无操作
CS2R	将特殊寄存器移动到寄存器
S2R	将特殊寄存器移动到寄存器
B2R	将屏障移至寄存器
BAR	屏障同步
R2B	将寄存器移动到屏障
VOTE	跨SIMD线程组投票

4.2. Volta指令集

Volta架构（计算能力7.0和7.2）具有以下指令集格式：

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

有效的目标位置和源位置包括：

寄存器RX
用于特殊系统控制寄存器的SRX
PX 用于谓词寄存器
c[X][Y] 用于常量内存

表5列出了适用于Volta GPU的有效指令集。

表5. Volta指令集
操作码	描述
浮点指令
FADD	FP32加法
FADD32I	浮点32位加法
FCHK	浮点数范围检查
FFMA32I	FP32 融合乘加运算
FFMA	FP32 融合乘加运算
FMNMX	FP32 最小值/最大值
FMUL	FP32乘法
FMUL32I	FP32乘法运算
FSEL	浮点选择
FSET	FP32比较并设置
FSETP	FP32比较并设置谓词
FSWZADD	FP32 交错加法
MUFU	FP32多功能运算
HADD2	FP16加法
HADD2_32I	FP16加法
HFMA2	FP16融合乘加运算
HFMA2_32I	FP16 融合乘加运算
HMMA	矩阵乘加运算
HMUL2	FP16乘法运算
HMUL2_32I	FP16乘法运算
HSET2	FP16比较并设置
HSETP2	FP16比较并设置谓词
DADD	双精度浮点数加法
DFMA	FP64 融合乘加运算
DMUL	双精度浮点数乘法
DSETP	FP64比较并设置谓词
整数指令
BMSK	位域掩码
BREV	位反转
FLO	查找前导1
IABS	整数绝对值
IADD	整数加法
IADD3	三输入整数加法
IADD32I	整数加法
IDP	整数点积与累加
IDP4A	整数点积累加运算
IMAD	整数乘加运算
IMMA	整数矩阵乘积累加
IMNMX	整数最小值/最大值
IMUL	整数乘法
IMUL32I	整数乘法
ISCADD	缩放整数加法
ISCADD32I	缩放整数加法
ISETP	整数比较与设置谓词
LEA	加载有效地址
LOP	逻辑运算
LOP3	逻辑运算
LOP32I	逻辑运算
POPC	人口计数
SHF	漏斗位移
SHL	左移
SHR	右移
VABSDIFF	绝对差值
VABSDIFF4	绝对差值
转换说明
F2F	浮点数到浮点数转换
F2I	浮点数转整数转换
I2F	整数到浮点数转换
I2I	整数到整数转换
I2IP	整数到整数转换与打包
FRND	四舍五入取整
移动指令
MOV	移动
MOV32I	移动
PRMT	寄存器对置换
SEL	带谓词的选择源
SGXT	符号扩展
SHFL	线程束级寄存器洗牌
谓词指令
PLOP3	谓词逻辑运算
PSETP	组合谓词并设置谓词
P2R	将谓词寄存器移动到寄存器
R2P	将寄存器移动到谓词寄存器
加载/存储指令
LD	从通用内存加载
LDC	加载常量
LDG	从全局内存加载
LDL	在本地内存窗口内加载
LDS	共享内存窗口内的加载
ST	存储到通用内存
STG	存储到全局内存
STL	存储到本地内存
STS	存储到共享内存
MATCH	跨线程组匹配寄存器值
QSPC	查询空间
ATOM	通用内存原子操作
ATOMS	共享内存原子操作
ATOMG	全局内存原子操作
RED	通用内存上的归约操作
CCTL	缓存控制
CCTLL	缓存控制
ERRBAR	错误屏障
MEMBAR	内存屏障
CCTLT	纹理缓存控制
纹理指令
TEX	纹理获取
TLD	纹理加载
TLD4	纹理加载4
TMML	纹理MipMap层级
TXD	带导数的纹理获取
TXQ	纹理查询
表面指令
SUATOM	表面内存原子操作
SULD	表面负载
SURED	表面内存上的归约操作
SUST	表面存储
控制指令
BMOV	移动收敛屏障状态
BPT	断点/陷阱
BRA	相对分支
BREAK	突破指定的收敛屏障
BRX	相对间接分支
BSSY	屏障集收敛同步点
BSYNC	在收敛屏障上同步线程
CALL	调用函数
EXIT	退出程序
JMP	绝对跳转
JMX	绝对间接跳转
KILL	终止线程
NANOSLEEP	暂停执行
RET	从子程序返回
RPCMOV	PC寄存器移动
RTT	从陷阱返回
WARPSYNC	同步线程束中的线程
YIELD	产出控制
其他指令
B2R	将屏障移至寄存器
BAR	屏障同步
CS2R	将特殊寄存器移动到寄存器
DEPBAR	依赖屏障
GETLMEMBASE	获取本地内存基地址
LEPC	加载有效程序计数器
NOP	无操作
PMTRIG	性能监控触发器
R2B	将寄存器移动到屏障
S2R	将特殊寄存器移动到寄存器
SETCTAID	设置CTA ID
SETLMEMBASE	设置本地内存基地址
VOTE	跨SIMD线程组投票

4.3. 图灵指令集

图灵架构（计算能力7.5）具有以下指令集格式：

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

有效的目标位置和源位置包括：

寄存器RX
用于统一寄存器的URX
用于特殊系统控制寄存器的SRX
用于谓词寄存器的PX
c[X][Y] 用于常量内存

表6列出了Turing GPU支持的有效指令。

表6. 图灵指令集
操作码	描述
浮点指令
FADD	FP32加法
FADD32I	浮点32位加法
FCHK	浮点数范围检查
FFMA32I	FP32 融合乘加运算
FFMA	FP32 融合乘加运算
FMNMX	FP32 最小值/最大值
FMUL	FP32乘法
FMUL32I	FP32乘法运算
FSEL	浮点选择
FSET	FP32比较并设置
FSETP	FP32比较并设置谓词
FSWZADD	FP32 交错加法
MUFU	FP32多功能运算
HADD2	FP16加法
HADD2_32I	FP16加法
HFMA2	FP16融合乘加运算
HFMA2_32I	FP16 融合乘加运算
HMMA	矩阵乘加运算
HMUL2	FP16乘法运算
HMUL2_32I	FP16乘法运算
HSET2	FP16比较并设置
HSETP2	FP16比较并设置谓词
DADD	双精度浮点数加法
DFMA	FP64 融合乘加运算
DMUL	双精度浮点数乘法
DSETP	FP64比较并设置谓词
整数指令
BMMA	位矩阵乘法累加
BMSK	位域掩码
BREV	位反转
FLO	查找前导1
IABS	整数绝对值
IADD	整数加法
IADD3	三输入整数加法
IADD32I	整数加法
IDP	整数点积与累加
IDP4A	整数点积累加运算
IMAD	整数乘加运算
IMMA	整数矩阵乘积累加
IMNMX	整数最小值/最大值
IMUL	整数乘法
IMUL32I	整数乘法
ISCADD	缩放整数加法
ISCADD32I	缩放整数加法
ISETP	整数比较与设置谓词
LEA	加载有效地址
LOP	逻辑运算
LOP3	逻辑运算
LOP32I	逻辑运算
POPC	人口计数
SHF	漏斗位移
SHL	左移
SHR	右移
VABSDIFF	绝对差值
VABSDIFF4	绝对差值
转换说明
F2F	浮点数到浮点数转换
F2I	浮点数转整数转换
I2F	整数到浮点数转换
I2I	整数到整数转换
I2IP	整数到整数转换与打包
FRND	四舍五入取整
移动指令
MOV	移动
MOV32I	移动
MOVM	带转置或扩展的矩阵移动
PRMT	寄存器对置换
SEL	带谓词的选择源
SGXT	符号扩展
SHFL	线程束级寄存器洗牌
谓词指令
PLOP3	谓词逻辑运算
PSETP	组合谓词并设置谓词
P2R	将谓词寄存器移动到寄存器
R2P	将寄存器移动到谓词寄存器
加载/存储指令
LD	从通用内存加载
LDC	加载常量
LDG	从全局内存加载
LDL	在本地内存窗口内加载
LDS	共享内存窗口内的加载
LDSM	从共享内存加载矩阵并进行元素大小扩展
ST	存储到通用内存
STG	存储到全局内存
STL	存储到本地内存
STS	存储到共享内存
MATCH	跨线程组匹配寄存器值
QSPC	查询空间
ATOM	通用内存原子操作
ATOMS	共享内存原子操作
ATOMG	全局内存原子操作
RED	通用内存上的归约操作
CCTL	缓存控制
CCTLL	缓存控制
ERRBAR	错误屏障
MEMBAR	内存屏障
CCTLT	纹理缓存控制
统一数据路径指令
R2UR	从向量寄存器移动到统一寄存器
S2UR	将特殊寄存器移动到统一寄存器
UBMSK	统一位域掩码
UBREV	统一位反转
UCLEA	为常量加载有效地址
UFLO	统一查找前导1
UIADD3	统一整数加法
UIADD3.64	统一整数加法
UIMAD	统一整数乘法
UISETP	整数比较并设置统一谓词
ULDC	从常量内存加载到统一寄存器
ULEA	统一加载有效地址
ULOP	逻辑运算
ULOP3	逻辑运算
ULOP32I	逻辑运算
UMOV	统一移动
UP2UR	统一谓词到统一寄存器
UPLOP3	统一谓词逻辑运算
UPOPC	统一人口计数
UPRMT	统一字节置换
UPSETP	统一谓词逻辑运算
UR2UP	统一寄存器到统一谓词
USEL	统一选择
USGXT	统一符号扩展
USHF	均匀漏斗偏移
USHL	统一左移
USHR	统一右移
VOTEU	在SIMD线程组中进行投票，结果存入统一目标
纹理指令
TEX	纹理获取
TLD	纹理加载
TLD4	纹理加载4
TMML	纹理MipMap层级
TXD	带导数的纹理获取
TXQ	纹理查询
表面指令
SUATOM	表面内存原子操作
SULD	表面负载
SURED	表面内存上的归约操作
SUST	表面存储
控制指令
BMOV	移动收敛屏障状态
BPT	断点/陷阱
BRA	相对分支
BREAK	突破指定的收敛屏障
BRX	相对间接分支
BRXU	基于统一寄存器偏移的相对分支
BSSY	屏障集收敛同步点
BSYNC	在收敛屏障上同步线程
CALL	调用函数
EXIT	退出程序
JMP	绝对跳转
JMX	绝对间接跳转
JMXU	基于统一寄存器偏移量的绝对跳转
KILL	终止线程
NANOSLEEP	暂停执行
RET	从子程序返回
RPCMOV	PC寄存器移动
RTT	从陷阱返回
WARPSYNC	同步线程束中的线程
YIELD	产出控制
其他指令
B2R	将屏障移至寄存器
BAR	屏障同步
CS2R	将特殊寄存器移动到寄存器
DEPBAR	依赖屏障
GETLMEMBASE	获取本地内存基地址
LEPC	加载有效程序计数器
NOP	无操作
PMTRIG	性能监控触发器
R2B	将寄存器移动到屏障
S2R	将特殊寄存器移动到寄存器
SETCTAID	设置CTA ID
SETLMEMBASE	设置本地内存基地址
VOTE	跨SIMD线程组投票

4.4. NVIDIA安培GPU与Ada指令集架构

NVIDIA Ampere GPU和Ada架构（计算能力8.0和8.6）具有以下指令集格式：

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

有效的目标位置和源位置包括：

寄存器RX
用于统一寄存器的URX
用于特殊系统控制寄存器的SRX
用于谓词寄存器的PX
用于统一谓词寄存器的UPX
c[X][Y] 用于常量内存

表7列出了适用于NVIDIA安培架构和Ada GPU的有效指令集。

表7. NVIDIA安培GPU与Ada指令集
操作码	描述
浮点指令
FADD	FP32加法
FADD32I	浮点32位加法
FCHK	浮点数范围检查
FFMA32I	FP32 融合乘加运算
FFMA	FP32 融合乘加运算
FMNMX	FP32 最小值/最大值
FMUL	FP32乘法
FMUL32I	FP32乘法运算
FSEL	浮点选择
FSET	FP32比较并设置
FSETP	FP32比较并设置谓词
FSWZADD	FP32 交错加法
MUFU	FP32多功能运算
HADD2	FP16加法
HADD2_32I	FP16加法
HFMA2	FP16融合乘加运算
HFMA2_32I	FP16 融合乘加运算
HMMA	矩阵乘加运算
HMNMX2	FP16 最小值/最大值
HMUL2	FP16乘法运算
HMUL2_32I	FP16乘法运算
HSET2	FP16比较并设置
HSETP2	FP16比较并设置谓词
DADD	双精度浮点数加法
DFMA	FP64 融合乘加运算
DMMA	矩阵乘积累加运算
DMUL	双精度浮点数乘法
DSETP	FP64比较并设置谓词
整数指令
BMMA	位矩阵乘法累加
BMSK	位域掩码
BREV	位反转
FLO	查找前导1
IABS	整数绝对值
IADD	整数加法
IADD3	三输入整数加法
IADD32I	整数加法
IDP	整数点积与累加
IDP4A	整数点积累加运算
IMAD	整数乘加运算
IMMA	整数矩阵乘积累加
IMNMX	整数最小值/最大值
IMUL	整数乘法
IMUL32I	整数乘法
ISCADD	缩放整数加法
ISCADD32I	缩放整数加法
ISETP	整数比较与设置谓词
LEA	加载有效地址
LOP	逻辑运算
LOP3	逻辑运算
LOP32I	逻辑运算
POPC	人口计数
SHF	漏斗位移
SHL	左移
SHR	右移
VABSDIFF	绝对差值
VABSDIFF4	绝对差值
转换说明
F2F	浮点数到浮点数转换
F2I	浮点数转整数转换
I2F	整数到浮点数转换
I2I	整数到整数转换
I2IP	整数到整数转换与打包
I2FP	整数转FP32格式转换与打包
F2IP	FP32向下转换为整数并打包
FRND	四舍五入取整
移动指令
MOV	移动
MOV32I	移动
MOVM	带转置或扩展的矩阵移动
PRMT	寄存器对置换
SEL	带谓词的选择源
SGXT	符号扩展
SHFL	线程束级寄存器洗牌
谓词指令
PLOP3	谓词逻辑运算
PSETP	组合谓词并设置谓词
P2R	将谓词寄存器移动到寄存器
R2P	将寄存器移动到谓词寄存器
加载/存储指令
LD	从通用内存加载
LDC	加载常量
LDG	从全局内存加载
LDGDEPBAR	全局加载依赖屏障
LDGSTS	异步全局内存到共享内存拷贝
LDL	在本地内存窗口内加载
LDS	共享内存窗口内的加载
LDSM	从共享内存加载矩阵并进行元素大小扩展
ST	存储到通用内存
STG	存储到全局内存
STL	存储到本地内存
STS	存储到共享内存
MATCH	跨线程组匹配寄存器值
QSPC	查询空间
ATOM	通用内存原子操作
ATOMS	共享内存原子操作
ATOMG	全局内存原子操作
RED	通用内存上的归约操作
CCTL	缓存控制
CCTLL	缓存控制
ERRBAR	错误屏障
MEMBAR	内存屏障
CCTLT	纹理缓存控制
统一数据路径指令
R2UR	从向量寄存器移动到统一寄存器
REDUX	将向量寄存器规约为统一寄存器
S2UR	将特殊寄存器移动到统一寄存器
UBMSK	统一位域掩码
UBREV	统一位反转
UCLEA	加载常量的有效地址
UF2FP	统一FP32向下转换与打包
UFLO	统一查找前导1
UIADD3	统一整数加法
UIADD3.64	统一整数加法
UIMAD	统一整数乘法
UISETP	整数比较并设置统一谓词
ULDC	从常量内存加载到统一寄存器
ULEA	统一加载有效地址
ULOP	逻辑运算
ULOP3	逻辑运算
ULOP32I	逻辑运算
UMOV	统一移动
UP2UR	统一谓词到统一寄存器
UPLOP3	统一谓词逻辑运算
UPOPC	统一人口计数
UPRMT	统一字节置换
UPSETP	统一谓词逻辑运算
UR2UP	统一寄存器到统一谓词
USEL	统一选择
USGXT	统一符号扩展
USHF	均匀漏斗偏移
USHL	统一左移
USHR	统一右移
VOTEU	在SIMD线程组中进行投票，结果存入统一目标
纹理指令
TEX	纹理获取
TLD	纹理加载
TLD4	纹理加载4
TMML	纹理MipMap层级
TXD	带导数的纹理获取
TXQ	纹理查询
表面指令
SUATOM	表面内存原子操作
SULD	表面负载
SURED	表面内存上的归约操作
SUST	表面存储
控制指令
BMOV	移动收敛屏障状态
BPT	断点/陷阱
BRA	相对分支
BREAK	突破指定的收敛屏障
BRX	相对间接分支
BRXU	基于统一寄存器偏移的相对分支
BSSY	屏障集收敛同步点
BSYNC	在收敛屏障上同步线程
CALL	调用函数
EXIT	退出程序
JMP	绝对跳转
JMX	绝对间接跳转
JMXU	基于统一寄存器偏移量的绝对跳转
KILL	终止线程
NANOSLEEP	暂停执行
RET	从子程序返回
RPCMOV	PC寄存器移动
WARPSYNC	同步线程束中的线程
YIELD	产出控制
其他指令
B2R	将屏障移至寄存器
BAR	屏障同步
CS2R	将特殊寄存器移动到寄存器
DEPBAR	依赖屏障
GETLMEMBASE	获取本地内存基地址
LEPC	加载有效程序计数器
NOP	无操作
PMTRIG	性能监控触发器
S2R	将特殊寄存器移动到寄存器
SETCTAID	设置CTA ID
SETLMEMBASE	设置本地内存基地址
VOTE	跨SIMD线程组投票

4.5. Hopper指令集

Hopper架构（计算能力9.0）具有以下指令集格式：

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

有效的目标位置和源位置包括：

寄存器RX
用于统一寄存器的URX
用于特殊系统控制寄存器的SRX
用于谓词寄存器的PX
用于统一谓词寄存器的UPX
c[X][Y] 用于常量内存
内存描述符的desc[URX][RY]
用于全局内存描述符的gdesc[URX]

表8列出了Hopper GPU支持的有效指令。

表8. Hopper指令集
操作码	描述
浮点指令
FADD	FP32加法
FADD32I	浮点32位加法
FCHK	浮点数范围检查
FFMA32I	FP32 融合乘加运算
FFMA	FP32 融合乘加运算
FMNMX	FP32 最小值/最大值
FMUL	FP32乘法
FMUL32I	FP32乘法运算
FSEL	浮点选择
FSET	FP32比较并设置
FSETP	FP32比较并设置谓词
FSWZADD	FP32 交错加法
MUFU	FP32多功能运算
HADD2	FP16加法
HADD2_32I	FP16加法
HFMA2	FP16融合乘加运算
HFMA2_32I	FP16 融合乘加运算
HMMA	矩阵乘加运算
HMNMX2	FP16 最小值/最大值
HMUL2	FP16乘法运算
HMUL2_32I	FP16乘法运算
HSET2	FP16比较并设置
HSETP2	FP16比较并设置谓词
DADD	双精度浮点数加法
DFMA	FP64 融合乘加运算
DMMA	矩阵乘积累加运算
DMUL	双精度浮点数乘法
DSETP	FP64比较并设置谓词
整数指令
BMMA	位矩阵乘法累加
BMSK	位域掩码
BREV	位反转
FLO	查找前导1
IABS	整数绝对值
IADD	整数加法
IADD3	三输入整数加法
IADD32I	整数加法
IDP	整数点积与累加
IDP4A	整数点积累加运算
IMAD	整数乘加运算
IMMA	整数矩阵乘积累加
IMNMX	整数最小值/最大值
IMUL	整数乘法
IMUL32I	整数乘法
ISCADD	缩放整数加法
ISCADD32I	缩放整数加法
ISETP	整数比较与设置谓词
LEA	加载有效地址
LOP	逻辑运算
LOP3	逻辑运算
LOP32I	逻辑运算
POPC	人口计数
SHF	漏斗位移
SHL	左移
SHR	右移
VABSDIFF	绝对差值
VABSDIFF4	绝对差值
VHMNMX	SIMD FP16 三输入最小值/最大值
VIADD	SIMD整数加法
VIADDMNMX	SIMD整数加法与融合最小/最大值比较
VIMNMX	SIMD整数最小值/最大值
VIMNMX3	SIMD整数三输入最小值/最大值
转换说明
F2F	浮点数到浮点数转换
F2I	浮点数转整数转换
I2F	整数到浮点数转换
I2I	整数到整数转换
I2IP	整数到整数转换与打包
I2FP	整数转FP32格式转换与打包
F2IP	FP32向下转换为整数并打包
FRND	四舍五入取整
移动指令
MOV	移动
MOV32I	移动
MOVM	带转置或扩展的矩阵移动
PRMT	寄存器对置换
SEL	带谓词的选择源
SGXT	符号扩展
SHFL	线程束级寄存器洗牌
谓词指令
PLOP3	谓词逻辑运算
PSETP	组合谓词并设置谓词
P2R	将谓词寄存器移动到寄存器
R2P	将寄存器移动到谓词寄存器
加载/存储指令
FENCE	共享内存或全局内存的可见性保证
LD	从通用内存加载
LDC	加载常量
LDG	从全局内存加载
LDGDEPBAR	全局加载依赖屏障
LDGMC	减少负载
LDGSTS	异步全局内存到共享内存拷贝
LDL	在本地内存窗口内加载
LDS	共享内存窗口内的加载
LDSM	从共享内存加载矩阵并进行元素大小扩展
STSM	将矩阵存储到共享内存
ST	存储到通用内存
STG	存储到全局内存
STL	存储到本地内存
STS	存储到共享内存
STAS	带显式同步的异步存储到分布式共享内存
同步	同步单元
MATCH	跨线程组匹配寄存器值
QSPC	查询空间
ATOM	通用内存原子操作
ATOMS	共享内存原子操作
ATOMG	全局内存原子操作
REDAS	基于分布式共享内存的显式同步异步归约
REDG	通用内存上的归约操作
CCTL	缓存控制
CCTLL	缓存控制
ERRBAR	错误屏障
MEMBAR	内存屏障
CCTLT	纹理缓存控制
统一数据路径指令
R2UR	从向量寄存器移动到统一寄存器
REDUX	将向量寄存器规约为统一寄存器
S2UR	将特殊寄存器移动到统一寄存器
UBMSK	统一位域掩码
UBREV	统一位反转
UCGABAR_ARV	CGA屏障同步
UCGABAR_WAIT	CGA屏障同步
UCLEA	加载常量的有效地址
UF2FP	统一FP32向下转换与打包
UFLO	统一查找前导1
UIADD3	统一整数加法
UIADD3.64	统一整数加法
UIMAD	统一整数乘法
UISETP	整数比较并设置统一谓词
ULDC	从常量内存加载到统一寄存器
ULEA	统一加载有效地址
ULEPC	统一负载有效PC
ULOP	逻辑运算
ULOP3	逻辑运算
ULOP32I	逻辑运算
UMOV	统一移动
UP2UR	统一谓词到统一寄存器
UPLOP3	统一谓词逻辑运算
UPOPC	统一人口计数
UPRMT	统一字节置换
UPSETP	统一谓词逻辑运算
UR2UP	统一寄存器到统一谓词
USEL	统一选择
USETMAXREG	释放、解除分配并重新分配寄存器
USGXT	统一符号扩展
USHF	均匀漏斗偏移
USHL	统一左移
USHR	统一右移
VOTEU	在SIMD线程组中进行投票，结果存入统一目标
Warpgroup指令
BGMMA	跨线程束的位矩阵乘积累加
HGMMA	跨线程组执行矩阵乘加运算
IGMMA	跨线程组的整数矩阵乘加运算
QGMMA	跨线程组执行的FP8矩阵乘加运算
WARPGROUP	线程组同步
WARPGROUPSET	设置Warpgroup计数器
张量内存访问指令
UBLKCP	批量数据复制
UBLKPF	批量数据预取
UBLKRED	从共享内存批量复制数据并进行归约操作
UTMACCTL	TMA缓存控制
UTMACMDFLUSH	TMA命令刷新
UTMALDG	从全局内存加载张量到共享内存
UTMAPF	张量预取
UTMAREDG	从共享内存到全局内存的张量存储（带归约操作）
UTMASTG	张量存储从共享内存到全局内存
纹理指令
TEX	纹理获取
TLD	纹理加载
TLD4	纹理加载4
TMML	纹理MipMap层级
TXD	带导数的纹理获取
TXQ	纹理查询
表面指令
SUATOM	表面内存原子操作
SULD	表面负载
SURED	表面内存上的归约操作
SUST	表面存储
控制指令
ACQBULK	等待批量释放状态的Warp状态
BMOV	移动收敛屏障状态
BPT	断点/陷阱
BRA	相对分支
BREAK	突破指定的收敛屏障
BRX	相对间接分支
BRXU	基于统一寄存器偏移的相对分支
BSSY	屏障集收敛同步点
BSYNC	在收敛屏障上同步线程
CALL	调用函数
CGAERRBAR	CGA错误屏障
ELECT	选举领导者线程
ENDCOLLECTIVE	重置MCOLLECTIVE掩码
EXIT	退出程序
JMP	绝对跳转
JMX	绝对间接跳转
JMXU	基于统一寄存器偏移量的绝对跳转
KILL	终止线程
NANOSLEEP	暂停执行
PREEXIT	依赖任务启动提示
RET	从子程序返回
RPCMOV	PC寄存器移动
WARPSYNC	同步线程束中的线程
YIELD	产出控制
其他指令
B2R	将屏障移至寄存器
BAR	屏障同步
CS2R	将特殊寄存器移动到寄存器
DEPBAR	依赖屏障
GETLMEMBASE	获取本地内存基地址
LEPC	加载有效程序计数器
NOP	无操作
PMTRIG	性能监控触发器
S2R	将特殊寄存器移动到寄存器
SETCTAID	设置CTA ID
SETLMEMBASE	设置本地内存基地址
VOTE	跨SIMT线程组投票

4.6. Blackwell指令集

Blackwell架构（计算能力10.0和12.0）具有以下指令集格式：

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

有效的目标位置和源位置包括：

寄存器RX
用于统一寄存器的URX
用于特殊系统控制寄存器的SRX
PX 用于谓词寄存器
用于统一谓词寄存器的UPX
c[X][Y] 用于常量内存
内存描述符的desc[URX][RY]
全局内存描述符的gdesc[URX]
用于张量内存的tmem[URX]

表8列出了适用于Blackwell GPU的有效指令。

表8. Blackwell指令集
操作码	描述
浮点指令
FADD	FP32加法
FADD2	FP32加法
FADD32I	浮点32位加法
FCHK	浮点数范围检查
FFMA32I	FP32 融合乘加运算
FFMA	FP32 融合乘加运算
FFMA2	FP32 融合乘加运算
FHADD	FP32加法
FHFMA	FP32 融合乘加运算
FMNMX	FP32 最小值/最大值
FMNMX3	3输入浮点最小值/最大值
FMUL	FP32乘法
FMUL2	FP32乘法运算
FMUL32I	FP32乘法运算
FSEL	浮点选择
FSET	FP32比较并设置
FSETP	FP32比较并设置谓词
FSWZADD	FP32 交错加法
MUFU	FP32多功能运算
HADD2	FP16加法
HADD2_32I	FP16加法
HFMA2	FP16融合乘加运算
HFMA2_32I	FP16 融合乘加运算
HMMA	矩阵乘加运算
HMNMX2	FP16 最小值/最大值
HMUL2	FP16乘法运算
HMUL2_32I	FP16乘法运算
HSET2	FP16比较并设置
HSETP2	FP16比较并设置谓词
DADD	双精度浮点数加法
DFMA	FP64 融合乘加运算
DMMA	矩阵乘积累加运算
DMUL	双精度浮点数乘法
DSETP	FP64比较并设置谓词
OMMA	跨Warp的FP4矩阵乘积累加运算
QMMA	跨Warp的FP8矩阵乘积累加运算
整数指令
BMSK	位域掩码
BREV	位反转
FLO	查找前导1
IABS	整数绝对值
IADD	整数加法
IADD3	三输入整数加法
IADD32I	整数加法
IDP	整数点积与累加
IDP4A	整数点积累加运算
IMAD	整数乘加运算
IMMA	整数矩阵乘积累加
IMNMX	整数最小值/最大值
IMUL	整数乘法
IMUL32I	整数乘法
ISCADD	缩放整数加法
ISCADD32I	缩放整数加法
ISETP	整数比较与设置谓词
LEA	加载有效地址
LOP	逻辑运算
LOP3	逻辑运算
LOP32I	逻辑运算
POPC	人口计数
SHF	漏斗位移
SHL	左移
SHR	右移
VABSDIFF	绝对差值
VABSDIFF4	绝对差值
VHMNMX	SIMD FP16 三输入最小值/最大值
VIADD	SIMD整数加法
VIADDMNMX	SIMD整数加法与融合最小/最大值比较
VIMNMX	SIMD整数最小值/最大值
VIMNMX3	SIMD整数三输入最小值/最大值
转换说明
F2F	浮点数到浮点数转换
F2I	浮点数转整数转换
I2F	整数到浮点数转换
I2I	整数到整数转换
I2IP	整数到整数转换与打包
I2FP	整数转FP32格式转换与打包
F2IP	FP32向下转换为整数并打包
FRND	四舍五入取整
移动指令
MOV	移动
MOV32I	移动
MOVM	带转置或扩展的矩阵移动
PRMT	寄存器对置换
SEL	带谓词的选择源
SGXT	符号扩展
SHFL	线程束级寄存器洗牌
谓词指令
PLOP3	谓词逻辑运算
PSETP	组合谓词并设置谓词
P2R	将谓词寄存器移动到寄存器
R2P	将寄存器移动到谓词寄存器
加载/存储指令
FENCE	共享内存或全局内存的可见性保证
LD	从通用内存加载
LDC	加载常量
LDG	从全局内存加载
LDGDEPBAR	全局加载依赖屏障
LDGMC	减少负载
LDGSTS	异步全局内存到共享内存拷贝
LDL	在本地内存窗口内加载
LDS	共享内存窗口内的加载
LDSM	从共享内存加载矩阵并进行元素大小扩展
STSM	将矩阵存储到共享内存
ST	存储到通用内存
STG	存储到全局内存
STL	存储到本地内存
STS	存储到共享内存
STAS	带显式同步的异步存储到分布式共享内存
同步	同步单元
MATCH	跨线程组匹配寄存器值
QSPC	查询空间
ATOM	通用内存原子操作
ATOMS	共享内存原子操作
ATOMG	全局内存原子操作
REDAS	基于分布式共享内存的显式同步异步归约
REDAS	基于分布式共享内存的显式同步异步归约
REDG	通用内存上的归约操作
CCTL	缓存控制
CCTLL	缓存控制
ERRBAR	错误屏障
MEMBAR	内存屏障
CCTLT	纹理缓存控制
统一数据路径指令
CREDUX	将向量寄存器耦合归约为统一寄存器
CS2UR	将常量内存中的值加载到统一寄存器
LDCU	将常量内存中的值加载到统一寄存器
R2UR	从向量寄存器移动到统一寄存器
REDUX	将向量寄存器规约为统一寄存器
S2UR	将特殊寄存器移动到统一寄存器
UBMSK	统一位域掩码
UBREV	统一位反转
UCGABAR_ARV	CGA屏障同步
UCGABAR_WAIT	CGA屏障同步
UCLEA	加载常量的有效地址
UFADD	统一浮点32位加法
UF2F	统一浮点到浮点转换
UF2FP	统一FP32向下转换与打包
UF2I	统一浮点到整数转换
UF2IP	统一FP32向下转换为整数并打包
UFFMA	统一FP32融合乘加运算
UFLO	统一查找前导1
UFMNMX	统一浮点最小值/最大值
UFMUL	统一FP32乘法
UFRND	均匀舍入到整数
UFSEL	统一浮点选择
UFSET	统一浮点数比较并设置
UFSETP	统一浮点数比较并设置谓词
UI2F	统一整数到浮点数转换
UI2FP	统一整数到FP32转换与打包
UI2I	统一饱和整数到整数转换
UI2IP	统一双饱和整数到整数转换与打包
UIABS	统一整数绝对值
UIMNMX	均匀整数最小值/最大值
UIADD3	统一整数加法
UIADD3.64	统一整数加法
UIMAD	统一整数乘法
UISETP	统一整数比较并设置统一谓词
ULEA	统一加载有效地址
ULEPC	统一负载有效PC
ULOP	统一逻辑运算
ULOP3	统一逻辑运算
ULOP32I	统一逻辑运算
UMOV	统一移动
UP2UR	统一谓词到统一寄存器
UPLOP3	统一谓词逻辑运算
UPOPC	统一人口计数
UPRMT	统一字节置换
UPSETP	统一谓词逻辑运算
UR2UP	统一寄存器到统一谓词
USEL	统一选择
USETMAXREG	释放、解除分配并重新分配寄存器
USGXT	统一符号扩展
USHF	均匀漏斗偏移
USHL	统一左移
USHR	统一右移
UGETNEXTWORKID	统一获取下一个工作ID
UMEMSETS	初始化共享内存
UREDGR	带释放的全局内存统一归约
USTGR	统一存储到全局内存并释放
UVIADD	统一SIMD整数加法
UVIMNMX	统一SIMD整数最小值/最大值
UVIRTCOUNT	虚拟资源管理
VOTEU	在SIMD线程组中进行投票，结果存入统一目标
张量内存访问指令
UBLKCP	批量数据复制
UBLKPF	批量数据预取
UBLKRED	从共享内存批量复制数据并进行归约操作
UTMACCTL	TMA缓存控制
UTMACMDFLUSH	TMA命令刷新
UTMALDG	从全局内存加载张量到共享内存
UTMAPF	张量预取
UTMAREDG	从共享内存到全局内存的张量存储（带归约操作）
UTMASTG	张量存储从共享内存到全局内存
Tensor Core内存指令
LDT	从张量内存加载矩阵到寄存器文件
LDTM	从张量内存加载矩阵到寄存器文件
STT	将矩阵从寄存器文件存储到张量内存
STTM	将矩阵从寄存器文件存储到张量内存
UTCATOMSWS	对SW状态寄存器执行原子操作
UTCBAR	张量核心屏障
UTCCP	从共享内存到张量内存的异步数据拷贝
UTCHMMA	统一矩阵乘积累加
UTCIMMA	统一矩阵乘积累加
UTCOMMA	统一矩阵乘加运算
UTCQMMA	统一矩阵乘加运算
UTCSHIFT	在张量内存中移动元素
纹理指令
TEX	纹理获取
TLD	纹理加载
TLD4	纹理加载4
TMML	纹理MipMap层级
TXD	带导数的纹理获取
TXQ	纹理查询
表面指令
SUATOM	表面内存原子操作
SULD	表面负载
SURED	表面内存上的归约操作
SUST	表面存储
控制指令
ACQBULK	等待批量释放状态的Warp状态
ACQSHMINIT	等待共享内存初始化释放状态线程束状态
BMOV	移动收敛屏障状态
BPT	断点/陷阱
BRA	相对分支
BREAK	突破指定的收敛屏障
BRX	相对间接分支
BRXU	基于统一寄存器偏移的相对分支
BSSY	屏障集收敛同步点
BSYNC	在收敛屏障上同步线程
CALL	调用函数
CGAERRBAR	CGA错误屏障
ELECT	选举领导者线程
ENDCOLLECTIVE	重置MCOLLECTIVE掩码
EXIT	退出程序
JMP	绝对跳转
JMX	绝对间接跳转
JMXU	基于统一寄存器偏移量的绝对跳转
KILL	终止线程
NANOSLEEP	暂停执行
PREEXIT	依赖任务启动提示
RET	从子程序返回
RPCMOV	PC寄存器移动
WARPSYNC	同步线程束中的线程
YIELD	产出控制
其他指令
B2R	将屏障移至寄存器
BAR	屏障同步
CS2R	将特殊寄存器移动到寄存器
DEPBAR	依赖屏障
GETLMEMBASE	获取本地内存基地址
LEPC	加载有效程序计数器
NOP	无操作
PMTRIG	性能监控触发器
S2R	将特殊寄存器移动到寄存器
SETCTAID	设置CTA ID
SETLMEMBASE	设置本地内存基地址
VOTE	跨SIMT线程组投票

5. cu++filt

cu++filt 将经过CUDA C++混淆的低级标识符解码（还原）为用户可读的名称。对于每个输入的字母数字单词，cu++filt的输出要么是解码后的CUDA C++名称（如果该名称能解码为CUDA C++名称），要么就是原始名称本身。

5.1. 使用指南

cu++filt 接受一个或多个由字母、数字、下划线、美元符号或点号组成的单词，并尝试解析它们。基本用法如下：

cu++filt [options] <symbol(s)>

要对整个文件（如二进制文件）进行名称还原，可以将文件内容通过管道传递给cu++filt，例如使用以下命令：

nm <input file> | cu++filt

要解构函数名称但不打印其参数类型，请使用以下命令：

cu++filt -p <symbol(s)>

要从混淆符号中跳过前导下划线，请使用以下命令：

cu++filt -_ <symbol(s)>

以下是cu++filt的示例输出：

$ cu++filt _Z1fIiEbl
bool f<int>(long)

如输出所示，符号_Z1fIiEbl已成功还原。

要移除函数签名和参数中的所有类型，请使用 -p 选项：

$ cu++filt -p _Z1fIiEbl
f<int>

要从混淆符号中跳过前导下划线，请使用-_选项：

$ cu++filt -_ __Z1fIiEbl
bool f<int>(long)

要对整个文件进行名称还原，将文件内容通过管道传递给cu++filt：

$ nm test.sm_70.cubin | cu++filt
0000000000000000 t hello(char *)
0000000000000070 t hello(char *)::display()
0000000000000000 T hello(int *)

无法进行名称还原的符号将按原样打印回标准输出:

$ cu++filt _ZD2
_ZD2

可以从命令行中解析多个符号：

$ cu++filt _ZN6Scope15Func1Enez _Z3fooIiPFYneEiEvv _ZD2
Scope1::Func1(__int128, long double, ...)
void foo<int, __int128 (*)(long double), int>()
_ZD2

5.2. 命令行选项

表9列出了cu++filt支持的命令行选项，以及每个选项功能的描述。

表9. `cu++filt` 命令行选项
选项	描述
`-_`	去除下划线。在某些系统中，CUDA编译器会在每个名称前添加下划线。此选项用于移除开头的下划线。cu++filt是否默认移除下划线取决于目标平台。
`-p`	在解构函数名称时，不显示该函数参数的类型。
`-h`	打印cu++filt的选项摘要并退出。
`-v`	打印此工具的版本信息。

5.3. 库可用性

cu++filt 也可作为静态库（libcufilt）使用，可链接到现有项目中。以下接口描述了其用法：

char* __cu_demangle(const char *id, char *output_buffer, size_t *length, int *status)

该接口可在SDK中的“nv_decode.h”文件中找到。

输入参数

id 输入混淆字符串。

output_buffer Pointer to where the demangled buffer will be stored. This memory must be allocated with malloc. If output-buffer is NULL, memory will be malloc’d to store the demangled name and returned through the function return value. If the output-buffer is too small, it is expanded using realloc.

length 如果用户提供了预分配的内存，则必须指定输出缓冲区的大小。解构器(demangler)在需要重新分配大小时会用到此参数。如果length非空，解构后缓冲区的长度将被存入length中。

status *status 被设置为以下值之一：

0 - 反混淆操作成功

-1 - 发生内存分配失败

-2 - 不是一个有效的混淆ID

-3 - 发生输入验证失败（一个或多个参数无效）

返回值

指向以NUL结尾的还原名称起始位置的指针，如果还原失败则为NULL。调用者需负责使用free释放该内存。

注意: 该函数是线程安全的。

示例用法

#include <stdio.h>
#include <stdlib.h>
#include "nv_decode.h"

int main()
{
  int     status;
  const char *real_mangled_name="_ZN8clstmp01I5cls01E13clstmp01_mf01Ev";
  const char *fake_mangled_name="B@d_iDentiFier";

  char* realname = __cu_demangle(fake_mangled_name, 0, 0, &status);
  printf("fake_mangled_name:\t result => %s\t status => %d\n", realname, status);
  free(realname);

  size_t size = sizeof(char)*1000;
  realname = (char*)malloc(size);
  __cu_demangle(real_mangled_name, realname, &size, &status);
  printf("real_mangled_name:\t result => %s\t status => %d\n", realname, status);
  free(realname);

  return 0;
}

这将打印：

fake_mangled_name:   result => (null)     status => -2
real_mangled_name:   result => clstmp01<cls01>::clstmp01_mf01()   status => 0

6. nvprune

nvprune 会修剪主机对象文件和库，使其仅包含指定目标的设备代码。

6.1. 使用说明

nvprune 每次运行时接受一个输入文件，并生成一个新的输出文件。基本用法如下：

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

输入文件必须是一个可重定位的主机对象或静态库（而非主机可执行文件），输出文件将保持相同格式。

必须使用–arch或–generate-code选项来指定要保留的目标架构。文件中所有其他设备代码将被丢弃。目标架构可以是sm_NN（cubin）或compute_NN（ptx）。

例如，以下操作将修剪libcublas_static.a，使其仅包含sm_70 cubin，而非通常存在的所有目标：

nvprune -arch sm_70 libcublas_static.a -o libcublas_static70.a

请注意，这意味着libcublas_static70.a将无法在其他架构上运行，因此仅应在为单一架构构建时使用。

6.2. 命令行选项

表10列出了nvprune支持的命令行选项，以及每个选项功能的描述。每个选项都有长名称和短名称，可以互换使用。

表10. `nvprune` 命令行选项
选项（长格式）	选项（短格式）	描述
`--arch 架构名称>,...`	`-arch`	指定将保留在对象或库中的NVIDIA GPU架构名称。
`--generate-code`	`-gencode`	此选项的格式与nvcc的--generate-code选项相同，提供了一种指定应保留在对象或库中的多种架构的方法。仅使用'code'值作为匹配目标。此选项允许的关键字：'arch','code'。
`--no-relocatable-elf`	`-no-relocatable-elf`	不要保留任何可重定位的ELF文件。
`--output-file`	`-o`	指定输出文件的名称和位置。
`--help`	`-h`	打印此工具的帮助信息。
`--options-file <file>,...`	`-optf`	从指定文件包含命令行选项。
`--version`	`-V`	打印此工具的版本信息。

7. 附录

7.1. JSON格式

nvdisasm的输出是人类可读的文本，未针对机器处理进行格式化。任何使用nvdisasm输出的工具都必须解析人类可读的文本，这可能很慢，并且对文本的任何微小更改都可能破坏解析器。

基于JSON的格式提供了一种高效且可扩展的方法，用于从nvdisasm输出机器可读数据。使用-json选项可以生成符合以下JSON模式定义的JSON文档。

{
    "$id": "https://nvidia.com/cuda/cuda-binary-utilities/index.html#json-format",
    "description": "A JSON schema for NVIDIA CUDA disassembler. The $id attribute is not a real URL but a unique identifier for the schema",
    "$schema": "https://json-schema.org/draft/2020-12/schema",
    "title": "A JSON schema for NVIDIA CUDA disassembler",
    "version": "12-8-0",
    "type": "array",
    "minItems": 2,
    "prefixItems": [
        {
            "$ref": "#/$defs/metadata"
        },
        {
            "description": "A list of CUDA functions",
            "type": "array",
            "minItems": 1,
            "items": {
                "$ref": "#/$defs/function"
            }
        }
    ],
    "$defs": {
        "metadata": {
            "type": "object",
            "properties": {
                "ELF": {
                    "$ref": "#/$defs/elf-metadata"
                },
                "SM": {
                    "type": "object",
                    "properties": {
                        "version": {
                            "$ref": "#/$defs/sm-version"
                        }
                    },
                    "required": [
                        "version"
                    ]
                },
                "SchemaVersion": {
                    "$ref": "#/$defs/version"
                },
                "Producer": {
                    "type": "string",
                    "description": "Name and version of the CUDA disassembler tool",
                    "maxLength": 1024
                },
                "Description": {
                    "type": "string",
                    "description": "A description that may be empty",
                    "maxLength": 1024
                },
                ".note.nv.tkinfo": {
                    "$ref": "#/$defs/Elf64_NV_TKinfo"
                }
            },
            "required": [
                "ELF",
                "SM",
                "SchemaVersion",
                "Producer",
                "Description"
            ]
        },
        "elf-metadata": {
            "type": "object",
            "properties": {
                "layout-id": {
                    "description": "Indicates the layout of the ELF file, part of the ELF header flags. Undocumented enum",
                    "type": "integer"
                },
                "ei_osabi": {
                    "description": "Operating system/ABI identification",
                    "type": "integer"
                },
                "ei_abiversion": {
                    "description": "ABI version",
                    "type": "integer"
                }
            },
            "required": [
                "layout-id",
                "ei_osabi",
                "ei_abiversion"
            ]
        },
        "sm-version": {
            "type": "object",
            "properties": {
                "major": {
                    "type": "integer"
                },
                "minor": {
                    "type": "integer"
                }
            },
            "required": [
                "major",
                "minor"
            ]
        },
        "version": {
            "type": "object",
            "properties": {
                "major": {
                    "type": "integer"
                },
                "minor": {
                    "type": "integer"
                },
                "revision": {
                    "type": "integer"
                }
            },
            "required": [
                "major",
                "minor",
                "revision"
            ]
        },
        "sass-instruction-attribute": {
            "type": "object",
            "additionalProperties": {
                "type": "string"
            }
        },
        "sass-instruction": {
            "type": "object",
            "properties": {
                "predicate": {
                    "type": "string",
                    "description": "Instruction predicate"
                },
                "opcode": {
                    "type": "string",
                    "description": "The instruction opcode. May be empty to indicate a gap between non-contiguous instructions"
                },
                "operands": {
                    "type": "string",
                    "description": "Instruction operands separated by commas"
                },
                "extra": {
                    "type": "string",
                    "description": "Optional field"
                },
                "other-attributes": {
                    "type": "object",
                    "description": "Additional instruction attributes encoded as a map of string:string key-value pairs. Example: {'control-flow': 'True'}",
                    "items": {
                        "type": "string"
                    }
                },
                "other-flags": {
                    "type": "array",
                    "description": "Aditional instruction attributes encoded as a list strings",
                    "items": {
                        "type": "string"
                    }
                }
            },
            "required": [
                "opcode"
            ]
        },
        "function": {
            "type": "object",
            "properties": {
                "function-name": {
                    "type": "string"
                },
                "start": {
                    "type": "integer",
                    "description": "The function's start virtual address"
                },
                "length": {
                    "type": "integer",
                    "description": "The function's length in bytes"
                },
                "other-attributes": {
                    "type": "array",
                    "items": {
                        "type": "string"
                    }
                },
                "sass-instructions": {
                    "type": "array",
                    "items": {
                        "$ref": "#/$defs/sass-instruction"
                    }
                }
            },
            "required": [
                "function-name",
                "start",
                "length",
                "sass-instructions"
            ]
        },
        "Elf64_NV_TKinfo": {
            "type": "object",
            "properties": {
                "tki_toolkitVersion": {
                    "type": "integer"
                },
                "tki_objFname": {
                    "type": "string"
                },
                "tki_toolName": {
                    "type": "string"
                },
                "tki_toolVersion": {
                    "type": "string"
                },
                "tki_toolBranch": {
                    "type": "string"
                },
                "tki_toolOptions": {
                    "type": "string"
                }
            },
            "required": [
                "tki_toolkitVersion",
                "tki_objFname",
                "tki_toolName",
                "tki_toolVersion",
                "tki_toolBranch",
                "tki_toolOptions"
            ]
        }
    }
}

关于sass-instruction对象的说明

other-attributes 可能包含 "control-flow": "True" 键值对，用于表示控制流指令。
第n个（从0开始）SASS指令的地址可以通过start + n * 指令大小计算得出。在Volta及之后的架构中，指令大小为16字节，而在之前的架构中为8字节。
JSON列表可能包含空的指令对象；这些对象会计入指令索引，因为它们表示非连续指令之间的间隔。
一个空的指令对象只有一个字段opcode，其值为空字符串："opcode": ""

以下是来自 nvdisasm -json 的示例输出

[
    // First element in the list: Metadata
    {
        // ELF Metadata
        "ELF": {
            "layout-id": 4,
            "ei_osabi": 51,
            "ei_abiversion": 7
        },
        // SASS code SM version: SM89 (16 bytes instructions)
        "SM": {
            "version": {
                "major": 8,
                "minor": 9
            }
        },
        "SchemaVersion": {
            "major": 12,
            "minor": 8,
            "revision": 0
        },
        // Release details of nvdisasm
        "Producer": "nvdisasm V12.8.14 Build r570_00.r12.8/compiler.35033008_0",
        "Description": ""
    },
    // Second element in the list: Functions
    [
        {
            "function-name": "_Z10exampleKernelv",
            // Function start address
            "start": 0,
            // Function length in bytes
            "length": 384,
            "other-attributes": [],
            // SASS instructions
            "sass-instructions": [
                {
                    // Instruction at 0x00
                    "opcode": "IMAD.MOV.U32",
                    "operands": "R1,RZ,RZ,c[0x0][0x28]"
                },
                {
                    // Instruction at 0x10 (16 bytes increment)
                    "opcode": "MOV",
                    "operands": "R0,0x0"
                },
                {
                    // Instruction at 0x20
                    "opcode": "IMAD.MOV.U32",
                    "operands": "R4,RZ,RZ,c[0x4][0x8]"
                },
                // [...]
                {
                    "opcode": "CALL.ABS.NOINC",
                    "operands": "R2",
                    // other-attributes is an optional that can indicate control flow instructions
                    "other-attributes": {
                        "control-flow": "True"
                    }
                },
                {
                    "opcode": "EXIT",
                    "other-attributes": {
                        "control-flow": "True"
                    }
                },
                {
                    "opcode": "NOP"
                }
            ]
        }
    ]
]

8. 通知

8.1. 通知

本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“NVIDIA”）对本文件所含信息的准确性或完整性不作任何明示或暗示的陈述或保证，并对其中可能存在的错误不承担任何责任。NVIDIA对于因使用此类信息而产生的后果、或因使用该信息导致的第三方专利或其他权利侵权概不负责。本文件不构成对开发、发布或交付任何材料（定义见下文）、代码或功能的承诺。

NVIDIA保留随时对本文件进行更正、修改、增强、改进以及任何其他变更的权利，恕不另行通知。

客户在下单前应获取最新的相关信息，并确认这些信息是最新且完整的。

除非NVIDIA与客户授权代表签署的单独销售协议中另有约定，否则NVIDIA产品的销售均以订单确认时提供的NVIDIA标准销售条款和条件为准（以下简称"销售条款"）。NVIDIA特此明确反对将任何客户通用条款适用于本文件所述NVIDIA产品的采购。本文件不直接或间接构成任何合同义务。

NVIDIA产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于那些可以合理预期NVIDIA产品故障或失灵会导致人身伤害、死亡、财产或环境损害的应用场景。NVIDIA对于在此类设备或应用中使用和/或包含NVIDIA产品不承担任何责任，因此客户需自行承担相关风险。

NVIDIA不声明或保证基于本文档的产品适用于任何特定用途。NVIDIA未必会对每个产品的所有参数进行测试。客户应全权负责评估和确定本文档所含信息的适用性，确保产品适合并满足客户计划的应用需求，并执行必要的应用测试以避免应用或产品出现故障。客户产品设计中的缺陷可能会影响NVIDIA产品的质量和可靠性，并可能导致超出本文档范围的其他或不同的条件和/或要求。对于任何因以下原因导致的故障、损坏、成本或问题，NVIDIA不承担任何责任：(i) 以违反本文档的任何方式使用NVIDIA产品或(ii) 客户产品设计。

本文档不授予任何NVIDIA专利权、版权或其他NVIDIA知识产权的明示或暗示许可。NVIDIA发布的关于第三方产品或服务的信息，不构成NVIDIA对这些产品或服务的使用许可或担保认可。使用此类信息可能需要获得第三方基于其专利或其他知识产权的许可，或需要获得NVIDIA基于其专利或其他知识产权的许可。

本文件中的信息仅可在获得NVIDIA事先书面批准、未经改动完整复制且完全符合所有适用的出口法律法规，并附带所有相关条件、限制和声明的情况下进行复制。

本文件及所有NVIDIA设计规格、参考板、文件、图纸、诊断工具、清单和其他文档（统称及单独称为"材料"）均以"现状"提供。NVIDIA不对材料作出任何明示或默示的保证，包括但不限于对不侵权、适销性和特定用途适用性的默示保证免责。在法律允许的最大范围内，NVIDIA不就因使用本文件导致的任何损害承担责任，包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论损害成因如何，也无论责任理论为何，即使NVIDIA已被告知发生此类损害的可能性。不论客户因任何原因可能遭受的任何损害，NVIDIA对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

8.2. OpenCL

OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。

8.3. 商标

NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。

1. 概述

1.1. 什么是CUDA二进制文件？

1.2. cuobjdump与nvdisasm的区别

1.3. 命令选项类型与表示法

2. cuobjdump

2.1. 使用指南

2.2. 命令行选项

3. nvdisasm

3.1. 使用指南

3.2. 命令行选项

4. 指令集参考

4.1. Maxwell与Pascal指令集

4.2. Volta指令集

4.3. 图灵指令集

4.4. NVIDIA安培GPU与Ada指令集架构

4.5. Hopper指令集

4.6. Blackwell指令集

5. cu++filt

5.1. 使用指南

5.2. 命令行选项

5.3. 库可用性

6. nvprune

6.1. 使用说明

6.2. 命令行选项

7. 附录

7.1. JSON格式

8. 通知

8.1. 通知

8.2. OpenCL

8.3. 商标

1.2. `cuobjdump`与`nvdisasm`的区别