PTX ISA 9.4：11. 指示符

.version  major.minor    // major, minor are integers

.version 3.1
.version 3.0
.version 2.3

.target stringlist         // comma separated list of target specifiers
string = { sm_120a, sm_120f, sm_120,          // sm_12x target architectures
           sm_121a, sm_121f, sm_121,          // sm_12x target architectures
           sm_110a, sm_110f, sm_110,          // sm_11x target architectures
           sm_100a, sm_100f, sm_100,          // sm_10x target architectures
           sm_101a, sm_101f, sm_101,          // sm_10x target architectures
           sm_103a, sm_103f, sm_103           // sm_10x target architectures
           sm_107a, sm_107f, sm_107           // sm_10x target architectures
           sm_90a, sm_90,                     // sm_9x target architectures
           sm_80, sm_86, sm_87, sm_88, sm_89, // sm_8x target architectures
           sm_70, sm_72, sm_75,               // sm_7x target architectures
           sm_60, sm_61, sm_62,               // sm_6x target architectures
           sm_50, sm_52, sm_53,               // sm_5x target architectures
           sm_30, sm_32, sm_35, sm_37,        // sm_3x target architectures
           sm_20,                             // sm_2x target architectures
           sm_10, sm_11, sm_12, sm_13,        // sm_1x target architectures
           texmode_unified, texmode_independent,   // texturing mode
           debug,                                  // platform option
           map_f64_to_f32 };                       // platform option

.target sm_10       // baseline target architecture
.target sm_13       // supports double-precision
.target sm_20, texmode_independent
.target sm_90       // baseline target architecture
.target sm_90a      // PTX using architecture-specific features
.target sm_100f     // PTX using family-specific features

.address_size  address-size
address-size = { 32, 64 };

// example directives
   .address_size 32       // addresses are 32 bit
   .address_size 64       // addresses are 64 bit
// example of directive placement within a module
.version 2.3
.target sm_20
.address_size 64
…
.entry foo () {
…
}

.entry kernel-name ( param-list )  kernel-body
.entry kernel-name  kernel-body

.entry cta_fft
.entry filter ( .param .b32 x, .param .b32 y, .param .b32 z )
{
    .reg .b32 %r<99>;
    ld.param.b32  %r1, [x];
    ld.param.b32  %r2, [y];
    ld.param.b32  %r3, [z];
    ...
}
.entry prefix_sum ( .param .align 4 .s32 pitch[8000] )
{
.reg .s32 %t;
ld.param::entry.s32  %t, [pitch];
…
}

.func {.attribute(attr-list)} fname {.noreturn} {.abi_preserve N} {.abi_preserve_control N} function-body
.func {.attribute(attr-list)} fname (param-list) {.noreturn} {.abi_preserve N} {.abi_preserve_control N} function-body
.func {.attribute(attr-list)} (ret-param) fname (param-list) {.abi_preserve N} {.abi_preserve_control N} function-body

.func (.reg .b32 rval) foo (.reg .b32 N, .reg .f64 dbl)
{
.reg .b32 localVar;
… use N, dbl;
other code;
mov.b32 rval,result;
ret;
}
…
call (fooval), foo, (val0, val1);  // return value in fooval
…
.func foo (.reg .b32 N, .reg .f64 dbl) .noreturn
{
.reg .b32 localVar;
… use N, dbl;
other code;
mov.b32 rval, result;
ret;
}
…
call foo, (val0, val1);
…
.func (.param .u32 rval) bar(.param .u32 N, .param .align 4 .b8 numbers[])
{
.reg .b32 input0, input1;
ld.param.b32   input0, [numbers + 0];
ld.param.b32   input1, [numbers + 4];
…
other code;
ret;
}
…
.param .u32 N;
.param .align 4 .b8 numbers[8];
st.param.u32    [N], 2;
st.param.b32    [numbers + 0], 5;
st.param.b32    [numbers + 4], 10;
call (rval), bar, (N, numbers);
…

.alias fAlias, fAliasee;

.visible .func foo(.param .u32 p) {
   ...
}
.visible .func bar(.param .u32 p);
.alias bar, foo;
.entry test()
{
      .param .u32 p;
      ...
      call foo, (p);       // call foo directly
       ...
       .param .u32 p;
       call bar, (p);        // call foo through alias
}
.entry filter ( .param .b32 x, .param .b32 y, .param .b32 z )
{
    .reg .b32 %r1, %r2, %r3;
    ld.param.b32  %r1, [x];
    ld.param.b32  %r2, [y];
    ld.param.b32  %r3, [z];
    ...
}

Label:   .branchtargets  list-of-labels ;

  .function foo () {
      .reg .u32 %r0;
      ...
      L1:
      ...
      L2:
      ...
      L3:
      ...
      ts: .branchtargets L1, L2, L3;
      @p brx.idx %r0, ts;
      ...
.function bar() {
.reg .u32 %r0;
…
N0:
…
N1:
…
N2:
…
N3:
…
N4:
…
ts: .branchtargets N<5>;
@p brx.idx %r0, ts;
…

Label:   .calltargets  list-of-functions ;

calltgt:  .calltargets  fastsin, fastcos;
...
@p   call  (%f1), %r0, (%x), calltgt;
...

 // no input or return parameters
label: .callprototype _ .noreturn {.abi_preserve N} {.abi_preserve_control N};
// input params, no return params
label: .callprototype _ (param-list) .noreturn {.abi_preserve N} {.abi_preserve_control N};
// no input params, // return params
label: .callprototype (ret-param) _ {.abi_preserve N} {.abi_preserve_control N};
// input, return parameters
label: .callprototype (ret-param) _ (param-list) {.abi_preserve N} {.abi_preserve_control N};

Fproto1: .callprototype  _ ;
Fproto2: .callprototype  _ (.param .f32 _);
Fproto3: .callprototype  (.param .u32 _) _ ;
Fproto4: .callprototype  (.param .u32 _) _ (.param .f32 _);
...
@p   call  (%val), %r0, (%f1), Fproto4;
...
// example of array parameter
Fproto5: .callprototype _ (.param .b8 _[12]);
Fproto6: .callprototype  _ (.param .f32 _) .noreturn;
…
@p   call  %r0, (%f1), Fproto6;
…
// example of .abi_preserve
Fproto7: .callprototype _ (.param .b32 _) .abi_preserve 10;
…
@p   call %r0, (%r1), Fproto7;
…

.maxnreg n

.entry foo .maxnreg 16 { ... }  // max regs per thread = 16

.maxntid nx
.maxntid nx, ny
.maxntid nx, ny, nz

.entry foo .maxntid 256       { ... }  // max threads = 256
.entry bar .maxntid 16,16,4   { ... }  // max threads = 1024

.reqntid nx
.reqntid nx, ny
.reqntid nx, ny, nz

.entry foo .reqntid 256       { ... }  // num threads = 256
.entry bar .reqntid 16,16,4   { ... }  // num threads = 1024

.minnctapersm ncta

.entry foo .maxntid 256 .minnctapersm 4 { ... }

.maxnctapersm ncta

.entry foo .maxntid 256 .maxnctapersm 4 { ... }

.noreturn

.func foo .noreturn { ... }

.pragma list-of-strings ;

.pragma "nounroll";    // disable unrolling in backend
// disable unrolling for current kernel
.entry foo .pragma “nounroll”; { … }

.abi_preserve N

.func bar() .abi_preserve 8
// Indirect call via call prototype
.func (.param .b32 out[30]) foo (.param .b32 in[30]) .abi_preserve 10 { … }
…
mov.b64 lpfoo, foo;
prot: .callprototype (.param .b32 out[30]) _ (.param .b32 in[30]) .abi_preserve 10;
call (out), lpfoo, (in), prot;

.abi_preserve_control N

.func foo() .abi_preserve_control 14
// Indirect call via call prototype
.func (.param .b32 out[30]) bar (.param .b32 in[30]) .abi_preserve_control 10 { … }
…
mov.b64 lpbar, bar;
prot: .callprototype (.param .b32 out[30]) _ (.param .b32 in[30]) .abi_preserve_control 10;
call (out), lpbar, (in), prot;

.minperctamemory size;
size = 32-bit unsigned integer (bytes)

.entry foo .minperctamemory 2048 {
    ...
}

@@DWARF dwarf-string
dwarf-string may have one of the
.byte   byte-list   // comma-separated hexadecimal byte values
.4byte  int32-list  // comma-separated hexadecimal integers in range [0..2^32-1]
.quad   int64-list  // comma-separated hexadecimal integers in range [0..2^64-1]
.4byte  label
.quad   label

@@DWARF .section .debug_pubnames, "", @progbits
@@DWARF .byte   0x2b, 0x00, 0x00, 0x00, 0x02, 0x00
@@DWARF .4byte  .debug_info
@@DWARF .4byte  0x000006b5, 0x00000364, 0x61395a5f, 0x5f736f63
@@DWARF .4byte  0x6e69616d, 0x63613031, 0x6150736f, 0x736d6172
@@DWARF .byte   0x00, 0x00, 0x00, 0x00, 0x00

.section section_name { dwarf-lines }
dwarf-lines have the following formats:
.b8    byte-list       // comma-separated list of integers
// in range [-128..255]
.b16   int16-list      // comma-separated list of integers
// in range [-2^15..2^16-1]
.b32   int32-list      // comma-separated list of integers
// in range [-2^31..2^32-1]
label:                 // Define label inside the debug section
.b64   int64-list      // comma-separated list of integers
// in range [-2^63..2^64-1]
.b32   label
.b64   label
.b32   label+imm       // a sum of label address plus a constant integer byte
// offset(signed, 32bit)
.b64   label+imm       // a sum of label address plus a constant integer byte
// offset(signed, 64bit)
.b32   label1-label2   // a difference in label addresses between labels in
// the same dwarf section (32bit)
.b64   label3-label4   // a difference in label addresses between labels in
// the same dwarf section (64bit)

.section .debug_pubnames
{
    .b32    LpubNames_end0-LpubNames_begin0
  LpubNames_begin0:
    .b8     0x2b, 0x00, 0x00, 0x00, 0x02, 0x00
    .b32    .debug_info
  info_label1:
    .b32    0x000006b5, 0x00000364, 0x61395a5f, 0x5f736f63
    .b32    0x6e69616d, 0x63613031, 0x6150736f, 0x736d6172
    .b8     0x00, 0x00, 0x00, 0x00, 0x00
  LpubNames_end0:
}
.section .debug_info
{
.b32 11430
.b8 2, 0
.b32 .debug_abbrev
.b8 8, 1, 108, 103, 101, 110, 102, 101, 58, 32, 69, 68, 71, 32, 52, 46, 49
.b8 0
.b32 3, 37, 176, -99
.b32 info_label1
.b32 .debug_loc+0x4
.b8 -11, 11, 112, 97
.b32 info_label1+12
.b64 -1
.b16 -5, -65535
}

.file file_index "filename" {, timestamp, file_size}

.file 1 "example.cu"
.file 2 "kernel.cu"
.file 1 "kernel.cu", 1339013327, 64118

.loc file_index line_number column_position
.loc file_index line_number column_position,function_name label {+ immediate }, inlined_at file_index2 line_number2 column_position2

    .loc 2 4237 0
L1:                        // line 4237, col 0 of file #2,
                           // inherited from mov
    mov.u32  %r1,%r2;      // line 4237, col 0 of file #2
    add.u32  %r2,%r1,%r3;  // line 4237, col 0 of file #2
...
L2:                        // line 4239, col 5 of file #2,
                           // inherited from sub
    .loc 2 4239 5
    sub.u32  %r2,%r1,%r3;  // line 4239, col 5 of file #2
    .loc 1 21 3
    .loc 1 9 3, function_name info_string0, inlined_at 1 21 3
    ld.global.u32   %r1, [gg]; // Function at line 9
    setp.lt.s32 %p1, %r1, 8;   // inlined at line 21
    .loc 1 27 3
    .loc 1 10 5, function_name info_string1, inlined_at 1 27 3
    .loc 1 15 3, function_name .debug_str+16, inlined_at 1 10 5
    setp.ne.s32 %p2, %r1, 18;
    @%p2 bra    BB2_3;
.section .debug_str {
info_string0:
 .b8 95  // _
 .b8 90  // z
 .b8 51  // 3
 .b8 102 // f
 .b8 111 // o
 .b8 111 // o
 .b8 118 // v
 .b8 0

info_string1:
 .b8 95  // _
 .b8 90  // z
 .b8 51  // 3
 .b8 98  // b
 .b8 97  // a
 .b8 114 // r
 .b8 118 // v
 .b8 0
 .b8 95  // _
 .b8 90  // z
 .b8 51  // 3
 .b8 99  // c
 .b8 97  // a
 .b8 114 // r
 .b8 118 // v
 .b8 0
}
plaintext
 code

.loc_intermediate file_index line_number column_position

.func (.param .b32 func_retval0) _Z17noinline_multiplyii(
    .param .b32 _Z17noinline_multiplyii_param_0,
    .param .b32 _Z17noinline_multiplyii_param_1
) {
    .reg .b32 %r<4>;
    .loc 1 71 0
    ld.param.b32  %r1, [_Z17noinline_multiplyii_param_0];
    ld.param.b32  %r2, [_Z17noinline_multiplyii_param_1];
    .loc 1 72 5
    .loc_intermediate 2 23 6               // tileIR_source.ir:23:6
    mul.lo.s32    %r3, %r2, %r1;
    .loc 1 73 5
    .loc_intermediate 2 24 0               // tileIR_source.ir:24:0
    add.lo.s32    %r3, %r3, %r1;
}
.file 1 "/path/to/vectorAdd.cu"
.file 2 "./tileIR_source.ir"

.nv_intermediate_source_section {
    .code_block {
        .ir_name:        ir-name-string
        .sourceFileName: file-index
        .source_begin
            source-text
        .source_end
    }
    [ .code_block { ... } ]*
}
ir-name-string = quoted string identifier (e.g., “TileIR”)
file-index     = unsigned integer matching a .file directive
source-text    = raw text of the intermediate IR source (heredoc)

.visible .entry _Z9vectorAddPf(
    .param .u64 .ptr .align 1 _Z9vectorAddPf_param_0
) {
    .loc 1 11 15
    .loc_intermediate 2 137 23
    ld.b64 %rd2, [%SP];
    .loc 1 12 15
    .loc_intermediate 3 24 23
    ld.b64 %rd3, [%SP];
}
.file 1 "./vectorAdd.cu"
.file 2 "./tileIR_source_A.ir"
.file 3 "./tileIR_source_B.ir"
.nv_intermediate_source_section {
    .code_block {
        .ir_name:        "TileIR"
        .sourceFileName: 2
        .source_begin
            __global__ void vectorAdd(float *v) { float a, b; *v = a + b; }
        .source_end
    }
    .code_block {
        .ir_name:        "TileIR"
        .sourceFileName: 3
        .source_begin
            __global__ void vectorSub(float *v) { float a, b; *v = a - b; }
        .source_end
    }
}

.extern identifier

.extern .global .b32 foo;  // foo is defined in another module

.visible identifier

.visible .global .b32 foo;  // foo will be externally visible

.weak identifier

.weak .func (.reg .b32 val) foo;  // foo will be externally visible

.common identifier

.common .global .u32 gbl;

.reqnctapercluster nx
.reqnctapercluster nx, ny
.reqnctapercluster nx, ny, nz

.entry foo .reqnctapercluster 2         { . . . }
.entry bar .reqnctapercluster 2, 2, 1   { . . . }
.entry ker .reqnctapercluster 3, 2      { . . . }

.explicitcluster

.entry foo .explicitcluster         { . . . }

.maxclusterrank n

.entry foo ..maxclusterrank 8         { . . . }

.blocksareclusters

.entry foo .reqntid 32, 32, 1 .reqnctapercluster 32, 32, 1 .blocksareclusters { ... }

.language lang-list   // comma separated list of strings or integers

.entry foo () .language "ptx"                    { ... }
.func bar ()  .language "nvvm", "ptx"            { ... }
.entry baz () .language "ptx", "cuda c++", 0x7   { ... }