PTX ISA 9.4：10. 特殊寄存器 • Sunky Blog

PTX 包含一组预定义的只读变量，它们以特殊寄存器的形式对外可见，通过 mov 或 cvt 指令进行访问。

特殊寄存器包括：

%tid
%ntid
%laneid
%warpid
%nwarpid
%ctaid
%nctaid
%smid
%nsmid
%gridid
%is_explicit_cluster
%clusterid
%nclusterid
%cluster_ctaid
%cluster_nctaid
%cluster_ctarank
%cluster_nctarank
%lanemask_eq, %lanemask_le, %lanemask_lt, %lanemask_ge, %lanemask_gt
%clock, %clock_hi, %clock64
%pm0, ..., %pm7
%pm0_64, ..., %pm7_64
%envreg0, ..., %envreg31
%globaltimer, %globaltimer_lo, %globaltimer_hi
%reserved_smem_offset_begin, %reserved_smem_offset_end, %reserved_smem_offset_cap, %reserved_smem_offset<2>
%total_smem_size
%aggr_smem_size
%dynamic_smem_size
%current_graph_exec
%perctamemoryoffset
%perctamemorysize

10.1. 特殊寄存器：%tid#

%tid

CTA 内的线程标识符。

语法（预定义）

.sreg .v4 .u32 %tid;                  // thread id vector
.sreg .u32 %tid.x, %tid.y, %tid.z;    // thread id components

说明

一个预定义的、只读的、按线程设置的特殊寄存器，用 CTA 内的线程标识符进行初始化。 %tid 特殊寄存器包含一个 1D、2D 或 3D 向量以匹配 CTA 形状；未使用维度上的 %tid 值为 0。第四个元素未使用，且始终返回零。各维度上的线程数量由预定义的特殊寄存器 %ntid 指定。

CTA 中的每个线程都有唯一的 %tid。

在每个 CTA 维度上，%tid 各分量的取值范围为 0 到 %ntid-1。

在 1D CTA 中 %tid.y == %tid.z == 0。在 2D CTA 中 %tid.z == 0。

保证以下关系成立：

<=  %tid.x <  %ntid.x
<=  %tid.y <  %ntid.y
<=  %tid.z <  %ntid.z

示例

mov.u32 %r1,%tid.x; // move tid.x to %rh

// legacy code accessing 16-bit components of %tid mov.u16 %rh,%tid.x; cvt.u32.u16 %r2,%tid.z; // zero-extend tid.z to %r2

10.2. 特殊寄存器：%ntid#

%ntid

每个 CTA 的线程 ID 数量。

语法（预定义）

.sreg .v4 .u32 %ntid;                   // CTA shape vector
.sreg .u32 %ntid.x, %ntid.y, %ntid.z;   // CTA dimensions

说明

一个预定义的、只读的特殊寄存器，用每个 CTA 维度上的线程 ID 数量进行初始化。 %ntid 特殊寄存器包含一个 3D 的 CTA 形状向量，保存各 CTA 维度。CTA 维度非零；第四个元素未使用，且始终返回零。CTA 中的线程总数为 (%ntid.x * %ntid.y * %ntid.z)。

%ntid.y == %ntid.z == 1 in 1D CTAs.
%ntid.z ==1 in 2D CTAs.

%ntid.{x,y,z} 的最大值如下：

目标架构	%ntid.x	%ntid.y	%ntid.z
`sm_1x`	512	512	64
`sm_20`, `sm_3x`, `sm_5x`, `sm_6x`, `sm_7x`, `sm_8x`, `sm_9x`, `sm_10x`, `sm_12x`	1024	1024	64

示例

// compute unified thread id for 2D CTA
mov.u32  %r0,%tid.x;
mov.u32  %h1,%tid.y;
mov.u32  %h2,%ntid.x;
mad.u32  %r0,%h1,%h2,%r0;
mov.u16  %rh,%ntid.x;      // legacy code

10.3. 特殊寄存器：%laneid#

%laneid

通道标识符。

语法（预定义）

.sreg .u32 %laneid;

说明

一个预定义的、只读的特殊寄存器，返回线程在 warp 内的通道。通道标识符的取值范围为零到 WARP_SZ-1。

示例

mov.u32  %r, %laneid;

10.4. 特殊寄存器：%warpid#

%warpid

warp 标识符。

语法（预定义）

.sreg .u32 %warpid;

说明

一个预定义的、只读的特殊寄存器，返回线程的 warp 标识符。warp 标识符在 CTA 内提供唯一的 warp 编号，但不在一个 grid 内的各 CTA 之间保持唯一。同一 warp 内所有线程的 warp 标识符相同。

注意，%warpid 返回的是读取瞬间线程所在的位置，但其值可能在执行过程中发生变化，例如在抢占之后线程被重新调度时。因此，如果 kernel 代码中需要一个虚拟 warp 索引，应当使用 %ctaid 和 %tid 来计算该值；%warpid 主要是为了让性能分析和诊断代码采样并记录诸如工作位置映射与负载分布之类的信息。

示例

mov.u32  %r, %warpid;

10.5. 特殊寄存器：%nwarpid#

%nwarpid

warp 标识符的数量。

语法（预定义）

.sreg .u32 %nwarpid;

说明

一个预定义的、只读的特殊寄存器，返回 warp 标识符的最大数量。

示例

mov.u32  %r, %nwarpid;

10.6. 特殊寄存器：%ctaid#

%ctaid

grid 内的 CTA 标识符。

语法（预定义）

.sreg .v4 .u32 %ctaid;                      // CTA id vector
.sreg .u32 %ctaid.x, %ctaid.y, %ctaid.z;    // CTA id components

说明

一个预定义的、只读的特殊寄存器，用 CTA grid 内的 CTA 标识符进行初始化。 %ctaid 特殊寄存器包含一个 1D、2D 或 3D 向量，取决于 CTA grid 的形状和秩。第四个元素未使用，且始终返回零。

保证以下关系成立：

<=  %ctaid.x <  %nctaid.x
<=  %ctaid.y <  %nctaid.y
<=  %ctaid.z <  %nctaid.z

示例

mov.u32  %r0,%ctaid.x;
mov.u16  %rh,%ctaid.y;   // legacy code

10.7. 特殊寄存器：%nctaid#

%nctaid

每个 grid 的 CTA ID 数量。

语法（预定义）

.sreg .v4 .u32 %nctaid                      // Grid shape vector
.sreg .u32 %nctaid.x,%nctaid.y,%nctaid.z;   // Grid dimensions

说明

一个预定义的、只读的特殊寄存器，用每个 grid 维度上的 CTA 数量进行初始化。 %nctaid 特殊寄存器包含一个 3D 的 grid 形状向量，每个元素的值至少为 1。第四个元素未使用，且始终返回零。

%nctaid.{x,y,z} 的最大值如下：

目标架构	%nctaid.x	%nctaid.y	%nctaid.z
`sm_1x`, `sm_20`	65535	65535	65535
`sm_3x`, `sm_5x`, `sm_6x`, `sm_7x`, `sm_8x`, `sm_9x`, `sm_10x`, `sm_12x`	2³¹ -1	65535	65535

示例

mov.u32  %r0,%nctaid.x;
mov.u16  %rh,%nctaid.x;     // legacy code

10.8. 特殊寄存器：%smid#

%smid

SM 标识符。

语法（预定义）

.sreg .u32 %smid;

说明

一个预定义的、只读的特殊寄存器，返回某个线程正在其上执行的处理器（SM）标识符。SM 标识符的取值范围为 0 到 %nsmid-1。不保证 SM 标识符的编号是连续的。

示例

mov.u32  %r, %smid;

10.9. 特殊寄存器：%nsmid#

%nsmid

SM 标识符的数量。

语法（预定义）

.sreg .u32 %nsmid;

说明

一个预定义的、只读的特殊寄存器，返回 SM 标识符的最大数量。由于不保证 SM 标识符的编号是连续的，%nsmid 可能大于设备中 SM 的物理数量。

示例

mov.u32  %r, %nsmid;

10.10. 特殊寄存器：%gridid#

%gridid

Grid 标识符。

语法（预定义）

.sreg .u64 %gridid;

说明

预定义、只读的特殊寄存器，以每 grid 的时间性 grid 标识符初始化。调试器使用 %gridid 来区分并发（小型）grid 内的 CTA 与 cluster。

在执行期间，可能会反复启动程序，每次启动都会开启一个 grid-of-CTAs。该变量为此上下文提供时间性 grid 启动编号。

对于 sm_1x 目标，%gridid 限于范围 [0..2¹⁶-1]。对于 sm_20， %gridid 限于范围 [0..2³²-1]。sm_30 支持完整的 64 位范围。

示例

mov.u64  %s, %gridid;  // 64-bit read of %gridid
mov.u32  %r, %gridid;  // legacy code with 32-bit %gridid

10.11. 特殊寄存器：%is_explicit_cluster#

%is_explicit_cluster

检查用户是否显式指定了 cluster 启动。

语法（预定义）

.sreg .pred %is_explicit_cluster;

说明

预定义、只读的特殊寄存器，以「用户是否显式指定了 cluster 启动」的谓词值初始化。

示例

.reg .pred p;
mov.pred  p, %is_explicit_cluster;

10.12. 特殊寄存器：%clusterid#

%clusterid

grid 内的 cluster 标识符。

语法（预定义）

.sreg .v4 .u32 %clusterid;
.sreg .u32 %clusterid.x, %clusterid.y, %clusterid.z;

说明

预定义、只读的特殊寄存器，以在每个维度上 grid 内的 cluster 标识符初始化。 grid 中的每个 cluster 都有唯一的标识符。

%clusterid 特殊寄存器包含一个 1D、2D 或 3D 向量，具体取决于 cluster 的形状与秩。第四个元素未使用，始终返回零。

保证以下关系成立：

<=  %clusterid.x <  %nclusterid.x
<=  %clusterid.y <  %nclusterid.y
<=  %clusterid.z <  %nclusterid.z

示例

.reg .b32 %r<2>;
.reg .v4 .b32 %rx;
mov.u32     %r0, %clusterid.x;
mov.u32     %r1, %clusterid.z;
mov.v4.u32  %rx, %clusterid;

10.13. 特殊寄存器：%nclusterid#

%nclusterid

每 grid 的 cluster 标识符数量。

语法（预定义）

.sreg .v4 .u32 %nclusterid;
.sreg .u32 %nclusterid.x, %nclusterid.y, %nclusterid.z;

说明

预定义、只读的特殊寄存器，以每个 grid 维度上的 cluster 数量初始化。

%nclusterid 特殊寄存器包含一个 3D grid 形状向量，以 cluster 为单位保存 grid 的各维度。第四个元素未使用，始终返回零。

有关 %nclusterid.{x,y,z} 最大值的详细信息，请参阅 Cuda Programming Guide。

示例

.reg .b32 %r<2>;
.reg .v4 .b32 %rx;
mov.u32     %r0, %nclusterid.x;
mov.u32     %r1, %nclusterid.z;
mov.v4.u32  %rx, %nclusterid;

10.14. 特殊寄存器：%cluster_ctaid#

%cluster_ctaid

cluster 内的 CTA 标识符。

语法（预定义）

.sreg .v4 .u32 %cluster_ctaid;
.sreg .u32 %cluster_ctaid.x, %cluster_ctaid.y, %cluster_ctaid.z;

说明

预定义、只读的特殊寄存器，以在每个维度上 cluster 内的 CTA 标识符初始化。 cluster 中的每个 CTA 都有唯一的 CTA 标识符。

%cluster_ctaid 特殊寄存器包含一个 1D、2D 或 3D 向量，具体取决于 cluster 的形状。第四个元素未使用，始终返回零。

保证以下关系成立：

<=  %cluster_ctaid.x <  %cluster_nctaid.x
<=  %cluster_ctaid.y <  %cluster_nctaid.y
<=  %cluster_ctaid.z <  %cluster_nctaid.z

示例

.reg .b32 %r<2>;
.reg .v4 .b32 %rx;
mov.u32     %r0, %cluster_ctaid.x;
mov.u32     %r1, %cluster_ctaid.z;
mov.v4.u32  %rx, %cluster_ctaid;

10.15. 特殊寄存器：%cluster_nctaid#

%cluster_nctaid

每 cluster 的 CTA 标识符数量。

语法（预定义）

.sreg .v4 .u32 %cluster_nctaid;
.sreg .u32 %cluster_nctaid.x, %cluster_nctaid.y, %cluster_nctaid.z;

说明

预定义、只读的特殊寄存器，以在每个维度上 cluster 内的 CTA 数量初始化。

%cluster_nctaid 特殊寄存器包含一个 3D grid 形状向量，以 CTA 为单位保存 cluster 的各维度。第四个元素未使用，始终返回零。

有关 %cluster_nctaid.{x,y,z} 最大值的详细信息，请参阅 Cuda Programming Guide。

示例

.reg .b32 %r<2>;
.reg .v4 .b32 %rx;
mov.u32     %r0, %cluster_nctaid.x;
mov.u32     %r1, %cluster_nctaid.z;
mov.v4.u32  %rx, %cluster_nctaid;

10.16. 特殊寄存器：%cluster_ctarank#

%cluster_ctarank

cluster 内跨所有维度的 CTA 标识符。

语法（预定义）

.sreg .u32 %cluster_ctarank;

说明

预定义、只读的特殊寄存器，以 cluster 内跨所有维度的 CTA 秩初始化。

保证以下关系成立：

0  <=  %cluster_ctarank <  %cluster_nctarank

示例

.reg .b32 %r;
mov.u32  %r, %cluster_ctarank;

10.17. 特殊寄存器：%cluster_nctarank#

%cluster_nctarank

cluster 内跨所有维度的 CTA 标识符数量。

语法（预定义）

.sreg .u32 %cluster_nctarank;

说明

预定义、只读的特殊寄存器，以 cluster 内跨所有维度的 CTA 数量初始化。

示例

.reg .b32 %r;
mov.u32  %r, %cluster_nctarank;

10.18. 特殊寄存器：%lanemask_eq#

%lanemask_eq

32 位掩码，在与线程在 warp 中的 lane 编号相等的位置上置位。

语法（预定义）

.sreg .u32 %lanemask_eq;

说明

预定义、只读的特殊寄存器，以一个 32 位掩码初始化，该掩码在与线程在 warp 中的 lane 编号相等的位置上置位。

示例

mov.u32     %r, %lanemask_eq;

10.19. 特殊寄存器：%lanemask_le#

%lanemask_le

32 位掩码，在小于或等于线程在 warp 中的 lane 编号的各位置上置位。

语法（预定义）

.sreg .u32 %lanemask_le;

说明

预定义、只读的特殊寄存器，以一个 32 位掩码初始化，该掩码在小于或等于线程在 warp 中的 lane 编号的各位置上置位。

示例

mov.u32     %r, %lanemask_le

10.20. 特殊寄存器：%lanemask_lt#

%lanemask_lt

32 位掩码，在小于线程在 warp 中的 lane 编号的各位置上置位。

语法（预定义）

.sreg .u32 %lanemask_lt;

说明

预定义、只读的特殊寄存器，以一个 32 位掩码初始化，该掩码在小于线程在 warp 中的 lane 编号的各位置上置位。

示例

mov.u32     %r, %lanemask_lt;

10.21. 特殊寄存器：%lanemask_ge#

%lanemask_ge

32 位掩码，在大于或等于线程在 warp 中的 lane 编号的各位置上置位。

语法（预定义）

.sreg .u32 %lanemask_ge;

说明

预定义、只读的特殊寄存器，以一个 32 位掩码初始化，该掩码在大于或等于线程在 warp 中的 lane 编号的各位置上置位。

示例

mov.u32     %r, %lanemask_ge;

10.22. 特殊寄存器：%lanemask_gt#

%lanemask_gt

32 位掩码，在大于线程在 warp 中的 lane 编号的各位置上置位。

语法（预定义）

.sreg .u32 %lanemask_gt;

说明

预定义、只读的特殊寄存器，以一个 32 位掩码初始化，该掩码在大于线程在 warp 中的 lane 编号的各位置上置位。

示例

mov.u32     %r, %lanemask_gt;

10.23. 特殊寄存器：%clock , %clock_hi#

%clock, %clock_hi

%clock: 预定义的只读 32 位无符号周期计数器。
%clock_hi: %clock64 特殊寄存器的高 32 位。

语法（预定义）

.sreg .u32 %clock;
.sreg .u32 %clock_hi;

说明

特殊寄存器 %clock 和 %clock_hi 都是 32 位无符号只读周期计数器，回绕时不产生任何提示。

示例

mov.u32 r1,%clock;
mov.u32 r2, %clock_hi;

10.24. 特殊寄存器：%clock64#

%clock64

预定义的只读 64 位无符号周期计数器。

语法（预定义）

.sreg .u64 %clock64;

说明

特殊寄存器 %clock64 是 64 位无符号只读周期计数器，回绕时不产生任何提示。

示例

mov.u64  r1,%clock64;

10.25. 特殊寄存器：%pm0 … %pm7#

%pm0 … %pm7

性能监视计数器。

语法（预定义）

.sreg .u32 %pm<8>;

说明

特殊寄存器 %pm0 … %pm7 是 32 位无符号只读性能监视计数器。它们的行为目前未定义。

示例

mov.u32  r1,%pm0;
mov.u32  r1,%pm7;

10.26. 特殊寄存器：%pm0_64 … %pm7_64#

%pm0_64 … %pm7_64

64 位性能监视计数器。

语法（预定义）

.sreg .u64 %pm0_64;
.sreg .u64 %pm1_64;
.sreg .u64 %pm2_64;
.sreg .u64 %pm3_64;
.sreg .u64 %pm4_64;
.sreg .u64 %pm5_64;
.sreg .u64 %pm6_64;
.sreg .u64 %pm7_64;

说明

特殊寄存器 %pm0_64 … %pm7_64 是 64 位无符号只读性能监视计数器。它们的行为目前未定义。

示例

mov.u32  r1,%pm0_64;
mov.u32  r1,%pm7_64;

10.27. 特殊寄存器：%envreg<32>#

%envreg<32>

由驱动程序定义的只读寄存器。

语法（预定义）

.sreg .b32 %envreg<32>;

说明

一组 32 个预定义的只读寄存器，用于记录 PTX 程序在 PTX 虚拟机之外的执行环境。这些寄存器由驱动程序在 kernel 启动前初始化，可以包含 CTA 范围或 grid 范围的取值。

这些寄存器的确切语义在驱动程序文档中定义。

示例

mov.b32      %r1,%envreg0;  // move envreg0 to %r1

10.28. 特殊寄存器：%globaltimer , %globaltimer_lo , %globaltimer_hi#

%globaltimer, %globaltimer_lo, %globaltimer_hi

%globaltimer: 预定义的 64 位全局纳秒定时器。
%globaltimer_lo: %globaltimer 的低 32 位。
%globaltimer_hi: %globaltimer 的高 32 位。

语法（预定义）

.sreg .u64 %globaltimer;
.sreg .u32 %globaltimer_lo, %globaltimer_hi;

说明

供 NVIDIA 工具使用的特殊寄存器。其行为取决于具体目标，在后续 GPU 上可能发生变化或被移除。当 JIT 编译到其他目标时，这些寄存器的取值不作规定。

示例

mov.u64  r1,%globaltimer;

10.29. 特殊寄存器：%reserved_smem_offset_begin , %reserved_smem_offset_end , %reserved_smem_offset_cap , %reserved_smem_offset_<2>#

%reserved_smem_offset_begin, %reserved_smem_offset_end, %reserved_smem_offset_cap, %reserved_smem_offset_<2>

%reserved_smem_offset_begin: 保留共享内存区域的起始位置。
%reserved_smem_offset_end: 保留共享内存区域的结束位置。
%reserved_smem_offset_cap: 保留共享内存区域的总大小。
%reserved_smem_offset_<2>: 保留共享内存区域内的偏移量。

语法（预定义）

.sreg .b32 %reserved_smem_offset_begin;
.sreg .b32 %reserved_smem_offset_end;
.sreg .b32 %reserved_smem_offset_cap;
.sreg .b32 %reserved_smem_offset_<2>;

说明

这些是预定义的只读特殊寄存器，其中包含为 NVIDIA 系统软件保留的共享内存区域的相关信息。该共享内存区域不对用户开放，从用户代码中访问这一区域会导致未定义行为。详情参见 CUDA Programming Guide。

示例

.reg .b32 %reg_begin, %reg_end, %reg_cap, %reg_offset0, %reg_offset1;

mov.b32 %reg_begin, %reserved_smem_offset_begin; mov.b32 %reg_end, %reserved_smem_offset_end; mov.b32 %reg_cap, %reserved_smem_offset_cap; mov.b32 %reg_offset0, %reserved_smem_offset_0; mov.b32 %reg_offset1, %reserved_smem_offset_1;

10.30. 特殊寄存器：%total_smem_size#

%total_smem_size

kernel 的某个 CTA 使用的共享内存总大小。

语法（预定义）

.sreg .u32 %total_smem_size;

说明

预定义的只读特殊寄存器，初始化时的值为启动时为该 kernel 的 CTA 所分配的共享内存总大小（含静态分配与动态分配，但不含为 NVIDIA 系统软件保留的共享内存）。

返回的大小以目标架构支持的共享内存分配单元大小的倍数表示。

分配单元取值如下：

目标架构	共享内存分配单元大小
`sm_2x`	128 字节
`sm_3x`, `sm_5x`, `sm_6x`, `sm_7x`	256 字节
`sm_8x`, `sm_9x`, `sm_10x`, `sm_12x`	128 字节

示例

mov.u32  %r, %total_smem_size;

10.31. 特殊寄存器：%aggr_smem_size#

%aggr_smem_size

kernel 的某个 CTA 使用的共享内存总大小。

语法（预定义）

.sreg .u32 %aggr_smem_size;

说明

预定义的只读特殊寄存器，初始化时的值为共享内存的总聚合大小，由两部分构成：启动时分配的用户共享内存大小（含静态分配与动态分配），以及为 NVIDIA 系统软件保留的共享内存区域大小。

示例

mov.u32  %r, %aggr_smem_size;

10.32. 特殊寄存器：%dynamic_smem_size#

%dynamic_smem_size

kernel 启动时动态分配的共享内存大小。

语法（预定义）

.sreg .u32 %dynamic_smem_size;

说明

kernel 启动时动态分配的共享内存大小。

预定义的只读特殊寄存器，初始化时的值为启动时为该 kernel 的 CTA 动态分配的共享内存大小。

示例

mov.u32  %r, %dynamic_smem_size;

10.33. 特殊寄存器：%current_graph_exec#

%current_graph_exec

当前正在执行的 CUDA 设备图的标识符。

语法（预定义）

.sreg .u64 %current_graph_exec;

说明

预定义的只读特殊寄存器，初始化时的标识符指向当前正在执行的 CUDA 设备图。如果正在执行的 kernel 不属于任何 CUDA 设备图，该寄存器为 0。

有关 CUDA 设备图的更多细节，参见 CUDA Programming Guide。

示例

mov.u64  r1, %current_graph_exec;

10.34. 特殊寄存器：%perctamemoryoffset#

%perctamemoryoffset

当前 CTA 在 per-CTA 全局内存中的基址偏移。

语法（预定义）

.sreg .u64 %perctamemoryoffset;

说明

预定义的只读 64 位特殊寄存器。返回当前 CTA 在 per-CTA 全局内存区域中所对应内存切片的 64 位地址。

只有在同时指定了 .minperctamemory 指示符的入口函数内部，使用该特殊寄存器才是有效的。

示例

mov.u64  %rd0, %perctamemoryoffset;

10.35. 特殊寄存器：%perctamemorysize#

%perctamemorysize

per-CTA 全局内存切片的大小，单位为字节。

语法（预定义）

.sreg .u32 %perctamemorysize;

说明

预定义的只读 32 位特殊寄存器，初始化时的值为 per-CTA 全局内存区域的大小（单位为字节）。该值即在 per-CTA 全局内存中相邻 CTA 区域之间的步长。

保证该值为 1024 字节的倍数。

只有在同时指定了 .minperctamemory 指示符的入口函数内部，使用该特殊寄存器才是有效的。

在不支持分配 per-CTA 全局内存的驱动程序上，该特殊寄存器的值为 0。在支持这种分配的驱动程序上，保证该值大于或等于 .minperctamemory 指示符所指定的大小。

示例

mov.u32  %r0, %perctamemorysize;