本节介绍由 ptxas 定义的 .pragma 字符串。
12.1. Pragma 字符串：“nounroll”#
"nounroll"
在优化后端编译器中禁用循环展开。
语法
.pragma "nounroll";
说明
"nounroll" pragma 是一条用于在优化后端编译器中禁用循环展开的
指示符。
"nounroll" pragma 允许出现在模块、入口函数和语句三个级别，
其含义如下：
- 模块作用域
-
禁用模块内所有循环的展开，包括位于
.pragma之前的循环。 - 入口函数作用域
-
禁用入口函数体内所有循环的展开。
- 语句级 pragma
-
禁用以当前块为循环头的那个循环的展开。
注意，若要在语句级达到预期效果，"nounroll" 指示符必须出现在目标循环的循环头
基本块中任何指令语句之前。循环头块定义为支配循环体内所有块、且为循环回边目标的块。出现在
循环头块之外的语句级 "nounroll" 指示符会被静默忽略。
示例
.entry foo (...)
.pragma "nounroll"; // do not unroll any loop in this function
{
...
}
.func bar (…)
{
…
L1_head:
.pragma “nounroll”; // do not unroll this loop
…
@p bra L1_end;
L1_body:
…
L1_continue:
bra L1_head;
L1_end:
…
}
12.2. Pragma 字符串：“used_bytes_mask”#
"used_bytes_mask"
用于指示 ld 操作数据中被使用字节的掩码。
语法
.pragma "used_bytes_mask mask";
说明
"used_bytes_mask" pragma 是一条根据所提供的掩码指定加载操作中
被使用字节的指示符。
"used_bytes_mask" pragma 需要写在某条加载指令之前，且需要获知该加载
操作所使用字节的信息。
若其后的指令不是加载指令，则该 pragma 会被忽略。
对于没有此 pragma 的加载指令，假定加载操作的所有字节都会被使用。
操作数 mask 是一个 32 位整数，其中置位的比特指示加载操作数据中被使用的
字节。
语义
Each bit in mask operand corresponds to a byte data where each set bit represents the used byte.
Most-significant bit corresponds to most-significant byte of data.
// For 4 bytes load with only lower 3 bytes used
.pragma “used_bytes_mask 0x7”;
ld.global.u32 %r0, [gbl]; // Higher 1 byte from %r0 is unused
// For vector load of 16 bytes with lower 12 bytes used
.pragma “used_bytes_mask 0xfff”;
ld.global.v4.u32 {%r0, %r1, %r2, %r3}, [gbl]; // %r3 unused
示例
.pragma "used_bytes_mask 0xfff";
ld.global.v4.u32 {%r0, %r1, %r2, %r3}, [gbl]; // Only lower 12 bytes used
12.3. Pragma 字符串：“enable_smem_spilling”#
"enable_smem_spilling"
为 CUDA kernel 启用共享内存溢出。
语法
.pragma "enable_smem_spilling";
说明
"enable_smem_spilling" pragma 是一条启用向共享内存进行寄存器溢出的指示符。
在溢出过程中，寄存器首先溢出到共享内存，当分配的共享内存写满后，任何额外的溢出都会被重定向到
局部内存。由于共享内存访问比局部内存更快，这可以通过降低访存延迟来提升性能。
"enable_smem_spilling" pragma 仅允许在函数作用域内使用。应用后，
它为指定函数启用共享内存溢出。
该 pragma 仅在某些场景和特定编译模式下使用有效。在以下情况下不允许使用该 pragma， 否则可能导致错误：
逐函数编译模式：例如分离编译、设备调试、含递归函数调用的整程序编译、可扩展整程序编译
使用动态分配共享内存的 kernel
使用
setmaxnreg指令的 kernel
注意
如果未显式指定启动界限，编译器将假设每个 CTA 可能的最大线程数，据此估算每个 CTA 分配的 共享内存及相应的溢出大小。但是，如果 kernel 启动时每个 CTA 的线程数少于估算值，每个 CTA 分配的共享内存可能超出编译器估计的大小，从而可能限制一个 SM 上可启动的 CTA 数量。因此， 在未指定启动界限的情况下使用该 pragma 可能导致性能回退。故建议仅在显式指定启动界限时使用 此 pragma。
示例
.entry foo (...)
{
...
.pragma "enable_smem_spilling"; // Enable shared memory spilling for this function
...
}
12.4. Pragma 字符串：“frequency”#
"frequency"
指定基本块的执行频率。
语法
.pragma "frequency n";
说明
"frequency" pragma 是一条指定基本块被执行线程执行次数的指示符。
优化编译器后端将此 pragma 视为用于优化的提示。
操作数 n 是一个指定执行频率的 64 位非负整数常量。
注意，若要使该 pragma 达到预期效果，应将其指定在基本块的开头。基本块定义为只有单一入口 和单一出口的直线指令序列。
示例
.entry foo (...)
{
.pragma "frequency 32";
...
}
12.5. Pragma 字符串：“mma_throughput”#
"mma_throughput"
为 mma 操作启用编译器的循环优化。
语法
.pragma "mma_throughput";
说明
"mma_throughput" pragma 是一条用于在编译器中为 mma 操作启用
循环优化的指示符。
该 pragma 仅允许在入口函数作用域使用。指定该 pragma 的入口函数将启用编译器优化。
示例
.entry test_kernel()
{
....
....
.pragma "mma_throughput";
....
....
}