PTX ISA 9.4：13. 发行说明

本节介绍 PTX ISA 及其实现的变更历史。第一节描述当前 PTX ISA 9.4 版发布中的 ISA 与实现变更，其余各节则提供回溯至 PTX ISA 2.0 版的历史版本变更记录。

表 71 展示了 PTX 版本历史。

表 72 展示了架构专用和系列专用 PTX 指令的发布历史。除 PTX 指令外，其他架构专用和系列专用的特性与构造将在以下小节中说明：

• Tensor Copy 指令的限制

• TensorCore 第 5 代矩阵形状的目标 ISA 说明

• TensorCore 第 5 代指令内存描述符的目标 ISA 说明

• tcgen05.mma 的块缩放

13.1. PTX ISA 9.4 版变更#

新增特性

PTX ISA 9.4 版引入以下新特性：

• 新增对 sm_107 目标架构的支持。

• 新增对目标架构 sm_107f 的支持，该架构支持系列专用特性。

• 新增对目标架构 sm_107a 的支持，该架构支持架构专用特性。

• 为 set 指令新增 .u8x4、.s8x4、.u16x2、.s16x2 指令类型支持。

• 为 cvt 指令新增 .pzo 限定符支持。

• 为 cvt 指令新增 .rz 舍入模式支持，可用于目标类型 .e4m3x2、.e5m2x2、.e2m3x2、.e3m2x2、.e2m1x2。

• 为 mbarrier.expect_tx、mbarrier.complete_tx、mbarrier.arrive、mbarrier.arrive_drop、cp.async.bulk、cp.async.bulk.tensor 与 tcgen05.commit 指令新增 .multicast::cluster::32b 限定符支持。

• 为 tcgen05.commit 指令新增 .sync_restrict::shared::read::mma::a 限定符支持。

• 为 cp.async.bulk.prefetch 与 cp.async.bulk.prefetch.tensor 指令新增 .level::eviction_priority 限定符支持。

• 新增对 applypriority.async.bulk 与 applypriority.async.bulk.tensor 指令的支持。

• 为 cp.async.bulk.tensor、cp.reduce.async.bulk.tensor、cp.async.bulk.prefetch.tensor 指令新增 .override::global_address 与 .override_attribute 限定符支持。

• 为 cp.async.bulk.tensor 与 cp.reduce.async.bulk.tensor 指令新增 .im2col_no_offs::w 限定符支持。

• 为 tcgen05.ld 与 tcgen05.ld.red 指令新增 .spcompress 限定符支持。

• 新增对 spcompress 与 spdecompress 指令的支持。

• 为 cvt 指令新增 .ue5m3x2 指令类型支持。

• 为 cvt 指令新增 .scaled::n1::ue8m0 限定符支持。

• 将 UE5M3 新增为 tcgen05.mma 与 tcgen05.mma.sp 指令的缩放类型。

• 为 tcgen05.mma、tcgen05.mma.sp、tcgen05.mma.ws 与 tcgen05.mma.ws.sp 指令新增 .kind::ti16 限定符支持。

• 新增对 add、sub、mul 与 fma 指令配合 .f16x2/.bf16x2/.f32x2 类型的支持。

• 新增对 add、sub、mul 与 fma 指令配合 FP8/FP6/FP4 x4 打包类型的支持。

• 为 tcgen05.alloc 与 tcgen05.dealloc 指令新增 .exclusive 限定符支持。

• 为 cp.async.bulk 与 cp.async.bulk.tensor 指令新增 .report_mechanism 限定符支持。

• 为 tcgen05.mma 指令新增 .decompress::lut::b 限定符支持。

• 为 tcgen05.mma 指令新增 .collector::b::* 限定符支持。

• 新增对 .loc_intermediate 指示符的支持，并新增用于中间 IR 源码位置映射的 .nv_intermediate_source_section 构造。

• 为 prefetch 指令新增 .valid_addr 与 .L1::32B 限定符支持。

• 为 atom、red、cp.reduce.async.bulk 与 multimem.cp.reduce.async.bulk 指令新增配合 .f32 类型使用的 .noftz 限定符支持。

• 为 ld 指令新增 .proxy::readonly 限定符支持。

• 新增对 .minperctamemory 指示符的支持，并新增用于 per-CTA 全局内存的 %perctamemoryoffset、%perctamemorysize 特殊寄存器。

• 为 ldmatrix 指令的 .m8n16 形状新增 .s8.s4 类型支持。

• 为 fence.proxy.async::generic 指令新增 .sync_restrict::shared::cluster::read 限定符支持。

• 为 fence.proxy 指令新增配合 .alias 限定符使用的 .acquire 与 .release 限定符支持。

• 新增对 fabric.try_get.async.tensor、fabric.try_put.async.tensor 与 fabric.try_red.async.tensor 指令的支持。

• 新增对 fabric.try_atom 指令的支持。

• 新增对 createpolicy.range.fabric 指令的支持。

• 为 fabric.try_get、fabric.try_put 与 fabric.try_red 指令新增 .level::cache_hint 限定符支持。

语义变更与澄清

• 无

13.2. PTX ISA 9.3 版变更#

新增特性

PTX ISA 9.3 版引入以下新特性：

• 新增对 mma_throughput pragma 的支持。

• 新增对 clmad 指令的支持。

• 为 mbarrier.test_wait 与 mbarrier.try_wait 指令新增 .phase_type::* 限定符支持。

• 为 mbarrier.test_wait 与 mbarrier.try_wait 指令新增 reportPredicate 与 reportValue 操作数支持。

• 为 mbarrier 指令新增 .layout 限定符支持，并新增对 mbarrier.check_layout 指令的支持。

• 新增对 multimem.st.async 与 multimem.red.async 指令的支持。

• 为 cp.async.bulk、cp.reduce.async.bulk、multimem.cp.async.bulk 与 multimem.cp.reduce.async.bulk 指令新增 .sem 与 .scope 限定符支持。

• 新增对 fabric.try_get、fabric.try_put、fabric.try_red、fabric.try_pullred、fabric.wait、fabric.submit 指令的支持。

• 为 fence 指令新增 fence.proxy.to_proxykind::from_proxykind_fabric 限定符支持。

• 新增对 .language 指示符的支持。

语义变更与澄清

• 澄清了 red.async 指令支持 sys 作用域，而在使用 .release 限定符时不支持 cluster 作用域。

13.3. PTX ISA 9.2 版变更#

新增特性

PTX ISA 9.2 版引入以下新特性：

• 为 add、sub、min、max、neg 指令新增 .u8x4 与 .s8x4 指令类型支持。

• 新增对 add.sat.{u16x2/s16x2/u32} 指令的支持。

• 为 st.async 指令新增 .b128 类型支持。

• 为 cp.async.bulk 指令新增 .ignore_oob 限定符支持。

• 为 cvt 指令新增 .bf16x2 目标类型支持，可用于源类型 .e4m3x2、.e5m2x2、.e3m2x2、.e2m3x2、.e2m1x2。

语义变更与澄清

• 对于 .atype 与 .btype 为 .e4m3/.e5m2、且 .dtype 为 .f32 的 wgmma.mma_async 指令，当前实现的累加精度高于半精度、但低于单精度。

13.4. PTX ISA 9.1 版变更#

新增特性

PTX ISA 9.1 版引入以下新特性：

• 为 ld 与 st 指令新增配合 .local 状态空间的 .volatile 限定符支持。

• 为 cvt 指令新增 .f16x2 与 .bf16x2 源类型支持，可用于目标类型 .e2m1x2、.e2m3x2、.e3m2x2、.e4m3x2、.e5m2x2。

• 为 mma/mma.sp 指令新增 .scale_vec::4X 支持，在 .kind::mxf4nvf4 下以 .ue8m0 作为 .stype。

• 为 cvt 指令新增 .s2f6x2 指令类型支持。

• 新增对 multimem.cp.async.bulk 与 multimem.cp.reduce.async.bulk 指令的支持。

语义变更与澄清

无。

13.5. PTX ISA 9.0 版变更#

新增特性

PTX ISA 9.0 版引入以下新特性：

• 新增对 sm_88 目标架构的支持。

• 新增对 sm_110 目标架构的支持。

• 新增对目标架构 sm_110f 的支持，该架构支持系列专用特性。

• 新增对目标架构 sm_110a 的支持，该架构支持架构专用特性。

• 新增对 pragma enable_smem_spilling 的支持，该 pragma 用于为一个函数启用共享内存溢出。

• 新增对 pragma frequency 的支持，该 pragma 用于指定基本块的执行频率。

• 新增对指示符 .blocksareclusters 的支持，该指示符用于指定 CUDA 线程块映射到 cluster。

• 扩展 st.bulk 指令的 size 操作数以支持 32 位长度。

• 新增对性能调优指示符 .abi_preserve 与 .abi_preserve_control 的支持，这两个指示符用于指定由函数调用方保留的数据寄存器与控制寄存器数量。

语义变更与澄清

• 同一 kernel 内的所有 tcgen05 指令（tcgen05.alloc、tcgen05.dealloc、tcgen05.relinquish_alloc_permit、tcgen05.cp、tcgen05.shift、tcgen05.mma、tcgen05.mma.sp、tcgen05.mma.ws, tcgen05.mma.ws.sp、tcgen05.commit）必须为 .cta_group 限定符指定相同的值。

13.6. PTX ISA 8.8 版变更#

新增特性

PTX ISA 8.8 版引入以下新特性：

• 新增对 sm_103 目标架构的支持。

• 新增对目标架构 sm_103a 的支持，该架构支持架构专用特性。

• 新增对 sm_121 目标架构的支持。

• 新增对目标架构 sm_121a 的支持，该架构支持架构专用特性。

• 引入了以 “f” 后缀表示的系列专用目标架构。用于系列专用目标的 PTX 与同一系列中的所有后续目标兼容。新增对 sm_100f、sm_101f、sm_103f、sm_120f、sm_121f 的支持。

• 扩展 min 与 max 指令以支持三个输入参数。

• 扩展 tcgen05.mma 指令，新增对 scale_vectorsize 限定符 .block16 与 .block32 以及 K 维度 96 的支持。

• 扩展 tensormap.replace 指令的 .field3 以支持 96B 重排模式。

• 新增对 tcgen05.ld.red 指令的支持。

• 扩展 ld、ld.global.nc 与 st 指令以支持 256b 加载/存储操作。

• 表 73 列出了在系列专用目标上受支持的特性：

  表 73 提升为系列专用架构的特性列表

  特性	受支持的目标
  .m16n8、.m16n16、.m8n16 形状以及 .b8 类型（用于 ldmatrix/stmatrix）	sm_100f、sm_101f、sm_120f
  tcgen05 的形状 .16x64b、.16x128b、.16x256b、.16x32bx2、.32x32b、.4x256b、.32x128b、.64x128b、.128x256b、.128x128b、.31x256b	sm_100f、sm_101f
  setmaxnreg	sm_100f、sm_101f、sm_120f
  .cta_group 修饰符	sm_100f、sm_101f
  cvt 配合 .e2m1x2、.e3m2x2、.e2m3x2、.ue8m0x2	sm_100f、sm_101f、sm_120f
  multimem 配合 .acc::f16 与 .e5m2、.e5m2x2、.e5m2x4、.e4m3、.e4m3x2、.e4m3x4 类型	sm_100f、sm_101f
  tensormap.replace	sm_100f、sm_101f、sm_120f
  tcgen05.ld.red	sm_101f、sm_103f
  tcgen05.ld/st/fence/wait/commit/cp/alloc/dealloc/relinquish_alloc_permit	sm_100f、sm_101f
  tcgen05.mma{.ws}{.sp}（kind::mxf4/kind::mxf4nvf4 的 .sp 除外）	sm_100f、sm_101f
  tcgen05 的 .kind::mxf4nvf4、.kind::mxf4、.kind::mxf8f6f4、.kind::f16、.kind::tf32、.kind::f8f6f4	sm_100f、sm_101f
  .ashift、.collector_usage 修饰符（用于 tcgen05）	sm_100f、sm_101f
  修饰符 .b8x16、.b6x16_p32、.b4x16_p64	sm_100f、sm_101f、sm_120f
  .block_scale 修饰符	sm_100f、sm_101f、sm_120f
  mma{.sp} 配合 .e3m2、.e2m3、.e2m1 类型以及 .kind、.block_scale、.scale_vec_size 修饰符（.sp 配合 mxf4/mxf4nvf4 的情况除外）	sm_120f
  .scale_vec::1X/2X/4X 修饰符	sm_120f
  .block16/.block32 修饰符（作为 scale_vec 的别名）	sm_100f、sm_101f
  .warpx2::02_13、.warpx2::01_23、.warpx4、.pack::16b、.unpack::16b 修饰符（用于 tcgen05）	sm_100f、sm_101f
  clusterlaunchcontrol.try_cancel 的 multicast::cluster::all	sm_100f、sm_101f、sm_120f
  .tile::scatter4、.tile::gather4、.im2col::w、.im2col::w::128	sm_100f、sm_101f
  redux.f32	sm_100f
  scale-input-d（用于 tcgen05）	sm_100f

语义变更与澄清

• 澄清了浮点转整数转换针对 NaN 输入的行为。

13.7. PTX ISA 8.7 版变更#

新增特性

PTX ISA 8.7 版引入以下新特性：

• 新增对 sm_120 目标架构的支持。

• 新增对目标架构 sm_120a 的支持，该架构支持架构专用特性。

• 扩展 tcgen05.mma 指令以支持 .kind::mxf4nvf4 与 .scale_vec::4X 限定符。

• 扩展 mma 指令以支持 .f16 类型累加器与 .m16n8k16 形状（配合 FP8 类型 .e4m3 与 .e5m2）。

• 扩展 cvt 指令以支持 .rs 舍入模式及目标类型 .e2m1x4、.e4m3x4、.e5m2x4、.e3m2x4、.e2m3x4。

• 扩展对 st.async 与 red.async 指令的支持，新增 .mmio、.release、.global 与 .scope 限定符支持。

• 扩展 tensormap.replace 指令，为 .elemtype 限定符新增对取值 13 至 15 的支持。

• 扩展 mma 与 mma.sp::ordered_metadata 指令，新增对类型 .e3m2/.e2m3/.e2m1 以及限定符 .kind、.block_scale、.scale_vec_size 的支持。

语义变更与澄清

• 澄清了在 .tile::gather4、.tile::scatter4 模式下，张量坐标需要指定为 {col_idx, row_idx0, row_idx1, row_idx2, row_idx3}，即 {x, y0, y1, y2, y3}，而不是 {x0, x1, x2, x3, y}。

• 更新了 tcgen05.mma 指令的指令描述符，以澄清为将来使用而保留的位。

13.8. PTX ISA 8.6 版变更#

新增特性

PTX ISA 8.6 版引入以下新特性：

• 新增对 sm_100 目标架构的支持。

• 新增对目标架构 sm_100a 的支持，该架构支持架构专用特性。

• 新增对 sm_101 目标架构的支持。

• 新增对目标架构 sm_101a 的支持，该架构支持架构专用特性。

• 扩展 cp.async.bulk 与 cp.async.bulk.tensor 指令，新增 .shared::cta 作为目标状态空间。

• 扩展 fence 指令以支持 .acquire 与 .release 限定符。

• 扩展 fence 与 fence.proxy 指令以支持 .sync_restrict 限定符。

• 扩展 ldmatrix 指令以支持 .m16n16、.m8n16 形状与 .b8 类型。

• 扩展 ldmatrix 指令以支持 .src_fmt、.dst_fmt 限定符。

• 扩展 stmatrix 指令以支持 .m16n8 形状与 .b8 类型。

• 新增对 clusterlaunchcontrol 指令的支持。

• 扩展 add、sub 与 fma 指令以支持混合精度浮点运算：以 .f32 作为目标操作数类型，以 .f16/.bf16 作为源操作数类型。

• 扩展 add、sub、mul 与 fma 指令以支持 .f32x2 类型。

• 扩展 cvt 指令的 .tf32 类型，以在 .rn/.rz 舍入模式下支持 .satfinite 限定符。

• 扩展 cp.async.bulk 指令以支持 .cp_mask 限定符与 byteMask 操作数。

• 扩展 multimem.ld_reduce 与 multimem.st 指令以支持 .e5m2、.e5m2x2、.e5m2x4、.e4m3、.e4m3x2 与 .e4m3x4 类型。

• 扩展 cvt 指令以支持与 .e2m1x2、.e3m2x2、.e2m3x2 及 .ue8m0x2 类型之间的转换。

• 扩展 cp.async.bulk.tensor 与 cp.async.bulk.prefetch.tensor 指令以支持新的 load_mode 限定符 .tile::scatter4 与 .tile::gather4。

• 扩展 tensormap.replace 指令以新增 .swizzle_atomicity 限定符支持，从而支持新的重排模式。

• 扩展 mbarrier.arrive、mbarrier.arrive_drop、.mbarrier.test_wait 与 .mbarrier.try_wait 指令以支持 .relaxed 限定符。

• 扩展 cp.async.bulk.tensor 与 cp.async.bulk.prefetch.tensor 指令以支持新的 load_mode 限定符 .im2col::w 与 .im2col::w::128。

• 扩展 cp.async.bulk.tensor 指令以支持新限定符 .cta_group。

• 新增对 st.bulk 指令的支持。

• 新增对 tcgen05 特性及相关指令的支持：tcgen05.alloc、tcgen05.dealloc、tcgen05.relinquish_alloc_permit、tcgen05.ld、tcgen05.st、tcgen05.wait、tcgen05.cp、tcgen05.shift、tcgen05.mma、tcgen05.mma.sp、tcgen05.mma.ws、tcgen05.mma.ws.sp、tcgen05.fence 与 tcgen05.commit。

• 扩展 redux.sync 指令以支持配合 .abs 与 .NaN 限定符使用的 .f32 类型。

语义变更与澄清

无。

13.9. PTX ISA 8.5 版变更#

新增特性

PTX ISA 8.5 版引入以下新特性：

• 新增对 mma.sp::ordered_metadata 指令的支持。

语义变更与澄清

• 稀疏性元数据（操作数 e）的取值 0b0000、0b0101、0b1010、0b1111 均为无效值，指令 mma.sp 对它们的使用会导致未定义行为。

13.10. PTX ISA 8.4 版变更#

新增特性

PTX ISA 8.4 版引入以下新特性：

• 使用 .b128 类型扩展 ld、st 和 atom 指令，以支持 .sys 作用域。

• 扩展整数 wgmma.mma_async 指令，分别支持将 .u8.s8 和 .s8.u8 用作 .atype 和 .btype。

• 扩展 mma 和 mma.sp 指令以支持 FP8 类型 .e4m3 和 .e5m2。

语义变更与澄清

无。

13.11. PTX ISA 8.3 版变更#

新增特性

PTX ISA 8.3 版引入以下新特性：

• 新增对 pragma used_bytes_mask 的支持，该 pragma 用于为加载操作指定已使用字节的掩码。

• 扩展 isspacep、cvta.to、ld 和 st 指令，使其接受 ::entry 和 ::func 子限定符，配合 .param 状态空间限定符使用。

• 为指令 ld、ld.global.nc、ldu、st、mov 和 atom 新增对 .b128 类型的支持。

• 新增对指令 tensormap.replace、tensormap.cp_fenceproxy 的支持，并新增对限定符 .to_proxykind::from_proxykind 在指令 fence.proxy 上的支持，以支持修改 tensor-map。

语义变更与澄清

无。

13.12. PTX ISA 8.2 版变更#

新增特性

PTX ISA 8.2 版引入以下新特性：

• 为 ld 和 st 指令新增对 .mmio 限定符的支持。

• 扩展 lop3 指令以允许谓词目标。

• 扩展 multimem.ld_reduce 指令以支持 .acc::f32 限定符，从而允许中间累加采用 .f32 精度。

• 扩展异步 Warpgroup 级矩阵乘累加运算 wgmma.mma_async 以支持 .sp 修饰符，该修饰符允许在输入矩阵 A 为稀疏时执行矩阵乘累加运算 。

语义变更与澄清

cp.async.bulk 和 cp.async.bulk.tensor 指令上的 .multicast::cluster 限定符已针对目标架构 sm_90a 进行了优化，在其他目标架构上性能可能大幅下降，因此建议将 .multicast::cluster 与 sm_90a 配合使用。

13.13. PTX ISA 8.1 版变更#

新增特性

PTX ISA 8.1 版引入以下新特性：

• 新增对 st.async 和 red.async 指令的支持，分别用于共享内存上的异步存储与 异步归约操作。

• 为半精度 fma 指令新增对 .oob 修饰符的支持。

• 为 cvt 指令在 .f16、.bf16、.tf32 格式上 新增对 .satfinite 饱和修饰符的支持。

• 将带 .e4m3/.e5m2 的 cvt 指令支持扩展至 sm_89。

• 扩展 atom 和 red 指令以支持向量类型。

• 新增对特殊寄存器 %aggr_smem_size 的支持。

• 以 64 位 min/max 操作扩展 sured 指令。

• 新增对增至 32764 字节的 kernel 参数大小的支持。

• 在内存一致性模型中新增对 multimem 地址的支持。

• 新增对 multimem.ld_reduce、multimem.st 和 multimem.red 指令的支持，用于在 multimem 地址上执行内存操作。

语义变更与澄清

无。

13.14. PTX ISA 8.0 版变更#

新增特性

PTX ISA 8.0 版引入以下新特性：

• 新增对支持架构特定特性的目标架构 sm_90a 的支持。

• 新增对异步 Warpgroup 级矩阵乘累加运算 wgmma 的支持。

• 以作用于大数据（包括张量数据）的批量操作扩展异步复制操作 。

• 引入打包整数类型 .u16x2 和 .s16x2。

• 扩展整数算术指令 add 以允许打包整数类型 .u16x2 和 .s16x2。

• 扩展整数算术指令 min 和 max 以允许打包整数类型 .u16x2 和 .s16x2，以及在 .s16x2 和 .s32 类型上的饱和修饰符 .relu 。

• 新增对特殊寄存器 %current_graph_exec 的支持，该寄存器标识当前正在执行的 CUDA 设备图。

• 新增对 elect.sync 指令的支持。

• 为函数和变量新增对 .unified 属性的支持。

• 新增对 setmaxnreg 指令的支持。

• 为 barrier.cluster 指令新增对 .sem 限定符的支持。

• 扩展 fence 指令，以允许使用 op_restrict 限定符进行指令操作码特定的同步。

• 新增对 .cluster 作用域的支持，可用于 mbarrier.arrive、mbarrier.arrive_drop、 mbarrier.test_wait 和 mbarrier.try_wait 操作。

• 在 mbarrier 对象上新增对事务计数操作的支持，通过 .expect_tx 和 .complete_tx 限定符指定。

语义变更与澄清

无。

13.15. PTX ISA 7.8 版变更#

新增特性

PTX ISA 7.8 版引入以下新特性：

• 新增对 sm_89 目标架构的支持。

• 新增对 sm_90 目标架构的支持。

• 扩展 bar 和 barrier 指令以接受可选作用域限定符 .cta。

• 以可选子限定符 ::cta 扩展 .shared 状态空间限定符。

• 新增对 movmatrix 指令的支持，该指令跨 warp 转置寄存器中的矩阵。

• 新增对 stmatrix 指令的支持，该指令将一个或多个矩阵存储到共享内存。

• 扩展 .f64 浮点类型 mma 运算，支持形状 .m16n8k4、.m16n8k8、 和 .m16n8k16。

• 扩展 add、sub、mul、set、setp、cvt、tanh、ex2、atom 和 red 指令以支持 bf16 备选浮点数据格式。

• 新增对备选浮点数据格式 .e4m3 和 .e5m2 的支持。

• 扩展 cvt 指令以转换 .e4m3 和 .e5m2 备选浮点数据格式。

• 新增对 griddepcontrol 指令的支持，作为控制相关 grid 执行的 通信机制。

• 扩展 mbarrier 指令以允许新的相位完成检查操作 try_wait。

• 新增对线程作用域 .cluster 的支持，它是一组协同线程阵列（CTA）。

• 扩展 fence/membar、ld、st、atom 和 red 指令以接受 .cluster 作用域。

• 新增对共享状态空间可见性的支持，可延伸至 cluster 内的所有线程。

• 以 ::cluster 子限定符扩展 .shared 状态空间限定符，实现 cluster 级的 共享内存可见性。

• 扩展 isspacep、cvta、ld、st、atom 和 red 指令以接受 配合 .shared 状态空间限定符的 ::cluster 子限定符。

• 新增对 mapa 指令的支持，将共享内存地址映射到 cluster 内不同 CTA 中的 相应地址。

• 新增对 getctarank 指令的支持，查询包含给定地址的 CTA 的 秩。

• 新增对新的栅栏同步指令 barrier.cluster 的支持。

• 扩展内存一致性模型以包含新的 cluster 作用域。

• 新增对与 cluster 信息相关的特殊寄存器 %is_explicit_cluster、 %clusterid、%nclusterid、%cluster_ctaid、%cluster_nctaid、%cluster_ctarank、 %cluster_nctarank 的支持。

• 新增对 cluster 维度指示符 .reqnctapercluster、.explicitcluster 和 .maxclusterrank 的支持。

语义变更与澄清

无。

13.16. PTX ISA 7.7 版变更#

新增特性

PTX ISA 7.7 版引入以下新特性：

• 扩展 isspacep 和 cvta 指令以纳入用于 kernel 函数参数的 .param 状态空间 。

语义变更与澄清

无。

13.17. PTX ISA 7.6 版变更#

新增特性

PTX ISA 7.6 版引入以下新特性：

• 支持 szext 指令，对指定值执行符号扩展或零扩展 。

• 支持 bmsk 指令，从指定位位置开始创建指定位宽的位掩码 。

• 支持特殊寄存器 %reserved_smem_offset_begin、%reserved_smem_offset_end、 %reserved_smem_offset_cap 和 %reserved_smem_offset<2>。

语义变更与澄清

无。

13.18. PTX ISA 7.5 版变更#

新增特性

PTX ISA 7.5 版引入以下新特性：

• 增强调试信息，以支持在 .section 调试指示符中使用标签差值和负值。

• 支持在 cp.async 指令上使用 ignore-src 操作数。

• 扩展内存一致性模型，引入以下新概念：

  • 将 内存代理 作为对不同内存访问方法的抽象标签。

  • 虚拟别名，即访问同一物理内存位置的不同内存地址。

• 支持新的 fence.proxy 和 membar.proxy 指令，以允许对通过虚拟别名执行的 内存访问进行同步。

语义变更与澄清

无。

13.19. PTX ISA 7.4 版变更#

新增特性

PTX ISA 7.4 版引入以下新特性：

• 支持 sm_87 目标架构。

• 支持 .level::eviction_priority 限定符，可在 ld、ld.global.nc、st 和 prefetch 指令上指定缓存驱逐优先级提示。

• 支持 .level::prefetch_size 限定符，可在 ld 和 cp.async 指令上指定数据预取提示。

• 支持 createpolicy 指令，可构造不同类型的缓存 驱逐策略。

• 支持 .level::cache_hint 限定符，可在 ld、ld.global.nc、st、atom、red 和 cp.async 指令上使用缓存驱逐策略。

• 支持对缓存数据执行 applypriority 和 discard 操作。

语义变更与澄清

无。

13.20. PTX ISA 7.3 版变更#

新增特性

PTX ISA 7.3 版引入以下新特性：

• 扩展初始值设定项中使用的 mask() 运算符，使其也支持整数常量表达式。

• 新增对栈操作指令的支持，这些指令允许使用 stacksave 和 stackrestore 指令操作栈，并使用 alloca 指令分配每线程栈。

语义变更与澄清

较早 PTX ISA 规范中未实现的 alloca 版本已在 PTX ISA 7.3 版中 由新的栈操作指令取代。

13.21. PTX ISA 7.2 版变更#

新增特性

PTX ISA 7.2 版引入以下新特性：

• 增强 .loc 指示符以表示内联函数信息。

• 新增支持在调试节中定义标签。

• 扩展 min 和 max 指令以支持 .xorsign 和 .abs 修饰符。

语义变更与澄清

无。

13.22. PTX ISA 7.1 版变更#

新增特性

PTX ISA 7.1 版引入以下新特性：

• 支持 sm_86 目标架构。

• 新增运算符 mask()，用于从初始值设定项中使用的变量地址提取特定字节 。

• 扩展 tex 和 tld4 指令以返回可选谓词，指示指定坐标处的数据是否 驻留在内存中。

• 扩展单位形式 wmma 和 mma 指令以支持 .and 运算。

• 扩展 mma 指令以支持 .sp 修饰符，该修饰符允许在输入矩阵 A 为稀疏时 执行矩阵乘累加运算。

• 扩展 mbarrier.test_wait 指令以测试特定相位奇偶校验的完成情况。

语义变更与澄清

无。

13.23. PTX ISA 7.0 版变更#

新增特性

PTX ISA 7.0 版引入以下新特性：

• 支持 sm_80 目标架构。

• 新增对异步复制指令的支持，这些指令允许将数据从一个状态空间异步复制到 另一个状态空间。

• 新增对 mbarrier 指令的支持，这些指令允许在内存中创建 mbarrier 对象， 并使用这些对象来同步线程以及由线程发起的异步复制操作。

• 新增对 redux.sync 指令的支持，该指令允许在 warp 内的线程间执行归约操作 。

• 新增对备选浮点数据格式 .bf16 和 .tf32 的支持。

• 扩展 wmma 指令以支持形状为 .m8n8k4 的 .f64 类型。

• 扩展 wmma 指令以支持 .bf16 数据格式。

• 扩展 wmma 指令以支持形状为 .m16n16k8 的 .tf32 数据格式。

• 扩展 mma 指令以支持形状为 .m8n8k4 的 .f64 类型。

• 扩展 mma 指令以支持 .bf16 和 .tf32 数据格式，形状为 .m16n8k8。

• 扩展 mma 指令以支持新形状 .m8n8k128、.m16n8k4、.m16n8k16、 .m16n8k32、.m16n8k64、.m16n8k128 和 .m16n8k256。

• 扩展 abs 和 neg 指令以支持 .bf16 和 .bf16x2 数据格式。

• 扩展 min 和 max 指令以支持 .NaN 修饰符及 .f16、.f16x2、 .bf16 和 .bf16x2 数据格式。

• 扩展 fma 指令以支持 .relu 饱和模式以及 .bf16 和 .bf16x2 数据格式。

• 扩展 cvt 指令以支持 .relu 饱和模式及 .f16、.f16x2、 .bf16、.bf16x2 和 .tf32 目标格式。

• 新增对 tanh 指令的支持，该指令计算双曲正切。

• 扩展 ex2 指令以支持 .f16 和 .f16x2 类型。

语义变更与澄清

无。

13.24. PTX ISA 6.5 版变更#

新增特性

PTX ISA 6.5 版引入以下新特性：

• 新增对半精度比较指令 set 的整数目标类型支持。

• 扩展 abs 指令以支持 .f16 和 .f16x2 类型。

• 新增对 cvt.pack 指令的支持，该指令允许转换两个整数值并将结果打包在一起。

• 为 mma 指令新增 .m16n8k8、.m8n8k16 和 .m8n8k32 形状。

• 新增对 ldmatrix 指令的支持，该指令为 mma 指令从共享内存加载一个或多个矩阵。

已移除特性

PTX ISA 6.5 版移除以下特性：

• 移除了在浮点 wmma.mma 指令上使用 .satfinite 限定符的支持。该支持自 PTX ISA 6.4 版起已弃用。

语义变更与澄清

无。

13.25. PTX ISA 6.4 版变更#

新增特性

PTX ISA 6.4 版引入以下新特性：

• 新增对 .noreturn 指示符的支持，可用于指示某个函数不返回其调用方函数。

• 新增对 mma 指令的支持，该指令允许执行矩阵乘加运算。

已弃用特性

PTX ISA 6.4 版弃用以下特性：

• 在浮点 wmma.mma 指令上使用 .satfinite 限定符的支持。

已移除特性

PTX ISA 6.4 版移除以下特性：

• 针对 .targetsm_70 及更高架构，移除了不带 .sync 限定符的 shfl 和 vote 指令支持。如 PTX ISA 6.2 版所述，该支持自 PTX ISA 6.0 版起已弃用。

语义变更与澄清

• 澄清了解析 .weak 符号的引用时，仅考虑同名的 .weak 或 .visible 符号，不考虑同名的局部符号。

• 澄清了在 cvt 指令中，仅当 .atype 或 .dtype 为 .f32 时才能指定修饰符 .ftz。

13.26. PTX ISA 6.3 版变更#

新增特性

PTX ISA 6.3 版引入以下新特性：

• 新增对 sm_75 目标架构的支持。

• 新增对一条新指令 nanosleep 的支持，该指令将线程挂起指定的时长。

• 新增对 .alias 指示符的支持，该指示符允许为函数符号定义别名。

• 扩展 atom 指令以执行 .f16 加法运算和 .cas.b16 运算。

• 扩展 red 指令以执行 .f16 加法运算。

• 扩展 wmma 指令以支持类型为 .s8、.u8、.s4、.u4、.b1 的被乘数矩阵以及类型为 .s32 的累加器矩阵。

语义变更与澄清

• 为所有 wmma 指令引入了必需的 .aligned 限定符。

• 规定了传递给 wmma.load 和 wmma.store 的基地址与步长参数所需的对齐。

• 澄清了 wmma 运算返回的 fragment 布局依赖于架构，并且在为不同链接兼容 SM 架构编译的函数之间传递 wmma fragment 可能无法按预期工作。

• 澄清了 {atom/red}.f16x2} 运算的原子性针对两个 .f16 元素分别保证，但不保证作为单个 32 位访问是原子的。

13.27. PTX ISA 6.2 版变更#

新增特性

PTX ISA 6.2 版引入以下新特性：

• 新增一条用于查询 warp 中活动线程的指令 activemask。

• 扩展原子与归约指令以执行 .f16x2 加法运算，需带必需的 .noftz 限定符。

已弃用特性

PTX ISA 6.2 版弃用以下特性：

• 不带 .sync 的 shfl 和 vote 指令的用法被追溯性地自 PTX ISA 6.0 版起弃用，该版本引入了实现独立线程调度的 sm_70 架构。

语义变更与澄清

• 澄清了 wmma 指令只能在已知 warp 中所有线程对条件的求值结果一致时，才可用于条件执行代码，否则行为未定义。

• 在内存一致性模型中，更新了道德强操作的定义，将栅栏从完全重叠的要求中排除，因为栅栏不访问内存。

13.28. PTX ISA 6.1 版变更#

新增特性

PTX ISA 6.1 版引入以下新特性：

• 新增对 sm_72 目标架构的支持。

• 新增对 wmma 指令中 32x8x16 和 8x32x16 新矩阵形状的支持。

语义变更与澄清

无。

13.29. PTX ISA 6.0 版变更#

新增特性

PTX ISA 6.0 版引入以下新特性：

• 新增对 sm_70 目标架构的支持。

• 规定了在 sm_70 及更高架构上运行的程序的内存一致性模型。

• 对内存指令的多项扩展，用于指定内存同步语义以及可观察到该同步的作用域。

• 新增用于矩阵运算的 wmma 指令，允许从内存加载矩阵、对其执行乘加运算并将结果存储到内存。

• 新增对新 barrier 指令的支持。

• 扩展 neg 指令以支持 .f16 和 .f16x2 类型。

• 新增一条 fns 指令，允许查找整数中第 n 个被置位的比特。

• 新增一条 bar.warp.sync 指令，允许同步 warp 中的线程。

• 扩展 vote 和 shfl 指令，加入 .sync 修饰符，在执行 vote 和 shfl 运算前分别等待指定线程。

• 新增一条 match.sync 指令，允许在 warp 中的线程间广播并比较一个值。

• 新增一条 brx.idx 指令，允许跳转到从候选目标列表中索引的标签。

• 新增对 .func 的无大小数组参数的支持，可用于实现可变参数函数。

• 新增对 dwarf-lines 中 .b16 整数类型的支持。

• 新增对使用 mov 指令取设备函数返回参数地址的支持。

语义变更与澄清

• 更新了 bar 指令的语义，指明执行线程会等待其 warp 中其他未退出的线程。

• 从规范中移除了 PTX 2.1 引入但从未实现的间接跳转支持。

• 从规范中移除了取标签地址、在初始化器中使用标签等从未实现的支持。

• 从规范中移除了从未实现的可变参数函数支持。

13.30. PTX ISA 5.0 版变更#

新增特性

PTX ISA 5.0 版引入以下新特性：

• 新增对 sm_60、sm_61、sm_62 目标架构的支持。

• 扩展原子与归约指令以执行双精度加法运算。

• 扩展原子与归约指令以指定 scope 修饰符。

• 新增 .common 指示符，允许链接包含同一符号不同大小声明的多个目标文件。

• 新增 dp4a 指令，允许执行带累加的 4 路点积运算。

• 新增 dp2a 指令，允许执行带累加的 2 路点积运算。

• 新增对特殊寄存器 %clock_hi 的支持。

语义变更与澄清

澄清了 ld 和 st 指令上缓存修饰符的语义，指明缓存运算仅被视为性能提示，不会改变程序的内存一致性行为。

澄清了 ld 和 st 指令上 volatile 运算的语义，指明优化编译器如何处理 volatile 运算。

13.31. PTX ISA 4.3 版变更#

新增特性

PTX ISA 4.3 版引入以下新特性：

• 新增 lop3 指令，允许对 3 个输入执行任意逻辑运算。

• 新增对扩展精度算术指令中 64 位计算的支持。

• 扩展 tex.grad 指令以支持 cube 和 acube 几何类型。

• 扩展 tld4 指令以支持 a2d、cube 和 acube 几何类型。

• 扩展 tex 和 tld4 指令以支持偏移向量和深度比较的可选操作数。

• 扩展 txq 指令以支持从特定 LOD 查询纹理字段。

语义变更与澄清

无。

13.32. PTX ISA 4.2 版变更#

新增特性

PTX ISA 4.2 版引入以下新特性：

• 新增对 sm_53 目标架构的支持。

• 新增针对 .f16 和 .f16x2 类型的算术、比较与纹理指令支持。

• 新增对曲面 memory_layout 字段的支持，以及用于查询该字段的 suq 指令支持。

语义变更与澄清

更新了 ABI 下参数传递的语义，指明用于参数传递的 ld.param 和 st.param 指令不能带谓词。

更新了 {atom/red}.add.f32 的语义，指明对全局内存的原子运算会将次正规输入和结果冲洗为保留符号的零；而对共享内存的原子运算会保留次正规输入和结果，不将其冲洗为零。

13.33. PTX ISA 4.1 版变更#

新增特性

PTX ISA 4.1 版引入以下新特性：

• 新增对 sm_37 和 sm_52 目标架构的支持。

• 新增对纹理的 array_size、num_mipmap_levels 和 num_samples 新字段的支持，以及用于查询这些字段的 txq 指令支持。

• 新增对曲面的 array_size 新字段的支持，以及用于查询该字段的 suq 指令支持。

• 新增对特殊寄存器 %total_smem_size 和 %dynamic_smem_size 的支持。

语义变更与澄清

无。

13.34. PTX ISA 4.0 版变更#

新增特性

PTX ISA 4.0 版引入以下新特性：

• 支持 sm_32 与 sm_50 目标架构。

• 支持 64 位性能计数器特殊寄存器 %pm0_64,..,%pm7_64。

• 新增 istypep 指令。

• 新增指令 rsqrt.approx.ftz.f64，用于计算值的平方根倒数的快速近似。

• 支持用于指定变量特殊属性的新指示符 .attribute。

• 支持 .managed 变量属性。

语义变更与澄清

vote 指令的语义已更新，以明确说明 warp 中的非活动线程在参与 vote.ballot.b32 时，其对应条目贡献 0。

13.35. PTX ISA 3.2 版变更#

新增特性

PTX ISA 3.2 版引入以下新特性：

• 纹理指令支持从多重采样纹理与多重采样数组纹理中读取。

• 扩展 .section 调试指示符，使其包含标签加立即数的表达式。

• 扩展 .file 指示符，使其包含时间戳与文件大小信息。

语义变更与澄清

vavrg2 与 vavrg4 指令的语义已更新，指明仅当 Va[i] + Vb[i] 非负时该指令才加 1，且加法结果右移 1 位（而不是除以 2）。

13.36. PTX ISA 3.1 版变更#

新增特性

PTX ISA 3.1 版引入以下新特性：

• 支持 sm_35 目标架构。

• 支持 CUDA 动态并行（CUDA Dynamic Parallelism），使 kernel 能够创建并同步新的工作。

• ld.global.nc，经由非一致性纹理缓存加载只读全局数据。

• 新增漏斗移位指令 shf。

• 扩展原子指令与归约指令，以执行 64 位 {and, or, xor} 运算以及 64 位整数 {min, max} 运算。

• 新增对 mipmaps 的支持。

• 新增对纹理与表面的间接访问支持。

• 将通用寻址的支持范围扩展到 .const 状态空间，并新增运算符 generic()，用于为初始化器中使用的 .global 或 .const 变量构造通用地址。

• 新增 .weak 指示符，允许链接包含同一符号声明的多个目标文件。

语义变更与澄清

PTX 3.1 重新定义了初始化器中全局变量的默认寻址方式，由通用地址改为全局状态空间中的偏移。遗留 PTX 代码被视为在初始化器中使用的每个全局变量都带有隐式的 generic() 运算符。PTX 3.1 代码应当要么在初始化器中包含显式的 generic() 运算符，要么使用 cvta.global 在运行时构造通用地址，要么使用 ld.global 从非通用地址加载。

mad.f32 指令要求在 sm_20 及更高的目标上使用舍入修饰符。但对于 PTX ISA 3.0 及更早版本，ptxas 并不强制这一要求，mad.f32 会静默地默认为 mad.rn.f32。对于 PTX ISA 3.1 版，ptxas 会生成一条警告并默认为 mad.rn.f32；在后续发布版本中，ptxas 将针对 PTX ISA 3.2 及更高版本强制执行这一要求。

13.37. PTX ISA 3.0 版变更#

新增特性

PTX ISA 3.0 版引入以下新特性：

• 支持 sm_30 目标架构。

• SIMD 视频指令。

• 新的 warp 洗牌指令。

• 指令 mad.cc 与 madc，用于高效的扩展精度整数乘法。

• 支持 3D 与数组几何的表面指令。

• 纹理指令支持从 cubemap 与 cubemap 数组纹理中读取。

• 平台选项 .target debug，用于声明某个 PTX 模块包含 DWARF 调试信息。

• pmevent.mask，用于触发多个性能监视事件。

• 性能监视计数器特殊寄存器 %pm4..%pm7。

语义变更与澄清

特殊寄存器 %gridid 已由 32 位扩展为 64 位。

PTX ISA 3.0 版在编译为应用二进制接口（ABI）时，弃用模块作用域的 .reg 与 .local 变量。在编译时不使用 ABI 的情况下，模块作用域的 .reg 与 .local 变量仍像以前一样受支持。编译包含模块作用域 .reg 或 .local 变量的遗留 PTX 代码（3.0 之前的 ISA 版本）时，编译器会静默地禁用 ABI 的使用。

shfl 指令的语义已更新，以明确说明对于 warp 内的非活动线程与被谓词关闭的线程，源操作数 a 的值是不可预测的。

PTX 模块不再允许重复的 .version 指示符。该特性此前从未实现，因此不存在语义变更。

未实现的指令 suld.p 与 sust.p.{u32,s32,f32} 已被移除。

13.38. PTX ISA 2.3 版变更#

新增特性

PTX 2.3 新增对纹理数组的支持。纹理数组特性支持访问由 1D 或 2D 纹理组成的数组：用一个整数索引纹理数组，随后使用一到两个单精度浮点坐标在所选的 1D 或 2D 纹理内进行寻址。

PTX 2.3 新增指示符 .address_size，用于指定地址的大小。

.const 与 .global 状态空间中的变量默认初始化为零。

语义变更与澄清

.maxntid 指示符的语义已更新，以与当前实现保持一致。具体而言，.maxntid 仅保证线程块中的线程总数不超过该最大值。此前的语义表明该最大值在各个维度上分别被强制执行，而事实并非如此。

位域提取与插入指令 BFE 与 BFI 现已指明 len 与 pos 操作数的取值限制在 0..255 范围内。

未实现的指令 {atom,red}.{min,max}.f32 已被移除。

13.39. PTX ISA 2.2 版变更#

新增特性

PTX 2.2 新增用于指定 kernel 参数属性的指示符；具体而言，新增了若干指示符，分别用于指定某个 kernel 参数是指针、指定该参数指向哪个状态空间，以及可选地指定该参数所指向内存的对齐要求。

PTX 2.2 在 .samplerref 不透明类型中新增了一个名为 force_unnormalized_coords 的字段。该字段用于独立纹理模式，以覆盖纹理头中的 normalized_coords 字段。之所以需要该字段，是为了支持 OpenCL 这类语言，它们在采样器头而非纹理头中表示归一化／未归一化坐标这一属性。

PTX 2.2 弃用显式常量存储体，并为 .const 状态空间支持一个较大的平坦地址空间。使用显式常量存储体的遗留 PTX 仍然受支持。

PTX 2.2 新增 tld4 指令，用于从构成给定纹理位置双线性插值覆盖范围的四个 texel 中加载某个分量（r、g、b 或 a）。该指令可用于在软件中计算精度更高的双线性插值结果，或用于执行带宽更高的纹理加载。

语义变更与澄清

无。

13.40. PTX ISA 2.1 版变更#

新增特性

PTX ISA 2.1 版为 sm_2x 目标支持底层的基于栈的 ABI。

sm_2x 目标已实现间接调用的支持。

新增指示符 .branchtargets 与 .calltargets，用于指定间接分支与间接函数调用的可能目标。新增 .callprototype 指示符，用于声明间接函数调用的类型签名。

.global 与 .const 变量的名称现在可以出现在变量初始化器中，以表示这些变量的地址。

新增一组 32 个驱动程序特定的执行环境特殊寄存器，命名为 %envreg0..%envreg31。

纹理与表面新增了通道数据类型与通道顺序字段，txq 与 suq 指令支持对这些字段进行查询。

指示符 .minnctapersm 取代了 .maxnctapersm 指示符。

新增指示符 .reqntid，允许指定确切的 CTA 维度。

新增指令 rcp.approx.ftz.f64，用于计算快速而粗略的近似倒数。

语义变更与澄清

若指定了 .minnctapersm 却未同时指定 .maxntid，则会输出一条警告。

13.41. PTX ISA 2.0 版变更#

新增特性

浮点扩展

本节介绍 PTX ISA 2.0 版针对 sm_20 目标的浮点变更。其目标是在尽可能实现 IEEE 754 兼容的同时，最大限度地保持与遗留 PTX ISA 1.x 代码及 sm_1x 目标的向后兼容性。

相对于 PTX ISA 1.x 版的变更如下：

• 对于 sm_20 目标，单精度指令默认支持次正规数。可使用 .ftz 修饰符来强制保持与 sm_1x 的兼容性。

• 对于 sm_20 目标，单精度 add、sub 与 mul 现已支持 .rm 与 .rp 舍入修饰符。

• 新增单精度融合乘加（fma）指令，支持符合 IEEE 754 的舍入修饰符并支持次正规数。fma.f32 指令还支持 .ftz 与 .sat 修饰符。fma.f32 要求 sm_20。mad.f32 指令已扩展出舍入修饰符，因此对于 sm_20 目标它与 fma.f32 含义相同。对于 sm_20 目标，fma.f32 与 mad.f32 都要求提供舍入修饰符。

• mad.f32 指令不带舍入的形式予以保留，以便编译器为 sm_1x 目标生成代码。为 sm_1x 编译的代码在 sm_20 设备上执行时，mad.f32 映射到 fma.rn.f32。

• 新增采用符合 IEEE 754 舍入的单精度与双精度 div、rcp 与 sqrt。这类形式通过是否使用舍入修饰符来标识，并要求 sm_20。

• 新增 testp 与 copysign 指令。

新增指令

新增统一加载指令 ldu。

表面指令支持额外的 .clamp 修饰符，即 .clamp 与 .zero。

sust 指令现已支持格式化表面存储。

新增统计前导零个数指令 clz。

新增查找前导非符号位指令 bfind。

新增位反转指令 brev。

新增位域提取与插入指令 bfe 与 bfi。

新增位计数指令 popc。

新增投票 ballot指令 vote.ballot.b32。

指令 {atom,red}.add.f32 已实现。

针对 sm_20 目标，指令 {atom,red}.shared 已扩展以处理 64 位数据类型。

新增系统级 membar 指令 membar.sys。

bar 指令的扩展如下：

• 新增 bar.arrive 指令。

• 新增 bar.red.popc.u32 与 bar.red.{and,or}.pred 指令。

• bar 现已支持可选的线程计数与寄存器操作数。

新增标量视频指令（包含 prmt）。

新增 isspacep 指令，用于查询某个通用地址是否落在指定状态空间的窗口之内。

新增 cvta 指令，用于在全局、局部与共享地址和通用地址之间进行相互转换。

其他新增特性

指令 ld、ldu、st、prefetch、prefetchu、isspacep、cvta、atom 与 red 现已支持通用寻址。

新增特殊寄存器 %nwarpid、%nsmid、%clock64、%lanemask_{eq,le,lt,ge,gt}。

为指令 ld、st、suld 与 sust 新增了缓存操作，例如用于将数据prefetching到内存层次结构中的指定层级。此外还新增了指令 prefetch 与 prefetchu。

指示符 .maxnctapersm 被弃用，并由 .minnctapersm 取代，以更贴合其行为与用途。

新增指示符 .section，用于取代 @@DWARF 语法，以通过 PTX 传递 DWARF 格式的调试信息。

新增指示符 .pragma nounroll，允许用户禁用循环展开。

语义变更与澄清

修复了 PTX ISA 1.4 及更早版本中 cvt.ftz 的勘误：当源类型或目标类型的大小为 64 位时，单精度次正规输入与结果未被冲洗为零。在 PTX ISA 1.5 及更高版本中，cvt.ftz（以及针对 .target sm_1x 的 cvt，其中隐含 .ftz）指令会在浮点指令类型的所有组合下，将单精度次正规输入与结果冲洗为保留符号的零。为保持与遗留 PTX 代码的兼容性，若 .version 为 1.4 或更早，则仅当源与目标类型的大小都不是 64 位时，单精度次正规输入与结果才被冲洗为保留符号的零。

特殊寄存器 %tid、%ntid、%ctaid 与 %nctaid 的各分量已由 16 位扩展为 32 位。这些寄存器现在的类型为 .v4.u32。

PTX ISA 1.5 版中曾错误地将独立纹理模式下可用的采样器数量列为三十二个；正确的数量是十六个。