SASS / PTX 指令速查：uniform datapath 与 warp 级原语

一、三层编译体系与可见性边界

在阅读任何反汇编之前，必须建立清晰的分层可见性模型。CUDA 源码到最终芯片执行经历了两级编译与三层表示：

+-------------------------------------------------------------------------+
|                              CUDA C++ 源码                              |
|   语义：Thread 抽象 (threadIdx, blockIdx)、指针、表达式、C++ 模板控制流   |
|   不可见：无物理寄存器、无 UR/UP 概念、无 SASS 指令排布与流水线延时       |
+-------------------------------------------------------------------------+
                                    │  nvcc (前端 Clang / EDG + NVVM 优化器)
                                    ▼
+-------------------------------------------------------------------------+
|                           PTX (虚拟中间 ISA)                            |
|   语义：无限虚拟寄存器 (%r<32>, %rd<64>, %p<8>)、强类型指令、显式同步原语  |
|   关键边界：PTX 完全没有 Uniform Register 类！所有寄存器均为每线程私有   |
+-------------------------------------------------------------------------+
                                    │  ptxas (汇编器 + 机器码生成器)
                                    ▼
+-------------------------------------------------------------------------+
|                           SASS (真实硬件机器码)                          |
|   语义：有限物理寄存器 (R0-R255, UR0-UR63, P0-P6, UP0-UP6)、发射控制码   |
|   关键边界：UR/UP 分配完全由 ptxas 基于 Divergence Analysis 静态推导生成  |
+-------------------------------------------------------------------------+

1. 核心事实：PTX 没有 Uniform Register

在 PTX 虚拟 ISA 中，寄存器声明只有 .reg .b32 %r<N>;、.reg .b64 %rd<N>;、.reg .pred %p<N>; 等，没有任何类似 .ureg 的标量寄存器类。从 PTX 的语法视角来看，所有变量都是每线程独立的（Per-thread）。

所有 Uniform Register（UR）以及 Uniform 谓词寄存器（UP）的分配与指令调度，完全发生在 ptxas 阶段。ptxas 会对 PTX 控制流与数据流进行静态分歧分析（Divergence Analysis）：

识别从根节点（如 blockIdx、gridDim、常数内存参数）出发、且在 Warp 内部 32 个线程中取值严格相同的数据流图；
将这些运算自动指派到硬件的 Uniform Datapath，并分配 UR0–UR63 存放中间值；
将统一的控制流条件指派给 Uniform 谓词 UP0–UP6，发射统一跳转或谓词执行。

2. 各层观察与分析工具链

观察目标	推荐命令	典型输出与关键观测点
观察 PTX	`nvcc -ptx -arch=sm_90 file.cu -o file.ptx`	检查高级指令是否生成 `bra.uni`、`redux.sync`、`shfl.sync`
从二进制提取 PTX	`cuobjdump -ptx file.o`	检查嵌入在 fatbin 中的 PTX 代码
反汇编 SASS	`cuobjdump -sass file.o`	查看完整机器码、控制码（Control Codes）与寄存器分配
源码对照反汇编	`nvdisasm -g -c file.o`	将 SASS 指令与 CUDA C++ 源码行号交织对齐
PTX 对照反汇编	`nvdisasm -gp -c file.o`	将 SASS 指令与 PTX 指令行号交织对齐，定位降低过程
寄存器与资源统计	`nvcc -Xptxas -v -arch=sm_90 -c file.cu`	观测 `Used X registers`（VGPR）与 `Used Y uniform registers`（UR）

二、寄存器与操作数速查

1. 寄存器分类与硬件语义

寄存器类	范围	硬件作用域	位宽	架构说明
`R0`–`R255`	256 个	Per-thread（向量寄存器）	32-bit	通用寄存器（VGPR）。每个 SM 总计 64K 个 32-bit 寄存器，按线程数分配。线程用量直接卡死 Warp Occupancy 档位。
`UR0`–`UR63`	64 个	Per-warp（标量寄存器）	32-bit	Turing（SM 7.5）引入的 Uniform Register。由同一个 Warp 内 32 个线程共享，不占每线程 VGPR 配额。
`RZ`	单一符号	Per-thread	32-bit	向量恒零寄存器（Zero Register）。读时恒为 0，写时作为丢弃目标的哑寄存器（Dummy Sink）。
`URZ`	单一符号	Per-warp	32-bit	Uniform 恒零寄存器。读时恒为 0，写时丢弃 uniform 运算结果。
`P0`–`P6` / `PT`	7 个 + 1 恒真	Per-thread	1-bit	向量条件谓词寄存器。`PT` 恒为真。用于 SIMT 分支选通与条件指令（如 `@P0 FADD ...`）。
`UP0`–`UP6` / `UPT`	7 个 + 1 恒真	Per-warp	1-bit	Uniform 谓词寄存器。`UPT` 恒为真。由 `UISETP` 产生，控制 Warp 级别无分歧分支或标量选通。

2. 常见特殊寄存器（SR_*）

特殊寄存器通过 S2R（读到向量寄存器）或 S2UR（读到 Uniform 寄存器）读取：

特殊寄存器	对应 CUDA C++ / PTX 内建变量	均匀性性质	典型读取指令示例	作用解释
`SR_TID.X` / `.Y` / `.Z`	`threadIdx.x / .y / .z`	Per-thread（发散）	`S2R R0, SR_TID.X`	线程局部三维索引
`SR_CTAID.X` / `.Y` / `.Z`	`blockIdx.x / .y / .z`	Warp-uniform	`S2UR UR4, SR_CTAID.X`	线程块三维坐标
`SR_LANEID`	`%laneid`（`threadIdx.x & 31`）	Per-thread（发散）	`S2R R11, SR_LANEID`	当前线程在 Warp 内的通道编号（0–31）
`SR_CgaCtaId`	`%cluster_ctarank`	Warp-uniform	`S2UR UR4, SR_CgaCtaId`	Hopper（SM 90+）Thread Block Cluster 内当前 CTA 编号
`SR_CLOCKLO` / `SR_CLOCKHI`	`%clock` / `%clock64`	动态递增	`CS2R.32 R15, SR_CLOCKLO`	SM 周期计数器（低 32 位 / 高 32 位）
`SR_GLOBALTIMERLO` / `HI`	`%globaltimer`	动态递增	`CS2R R2, SR_GLOBALTIMERLO`	全局纳秒单调时钟
`SR_VIRTUALSMID`	`%smid`	Warp-uniform	`S2UR UR8, SR_VIRTUALSMID`	当前执行所在的虚拟 SM 编号
`SR_VIRTCFG`	虚拟环境配置元数据	Warp-uniform	`S2UR UR9, SR_VIRTCFG`	Hopper / 容器虚拟化配置掩码

3. 常数内存（Constant Bank）常见偏移

SASS 中的 c[0x0][0x...] 表示 Constant Memory Bank 0，承载 Kernel 启动参数与运行时元数据。在 64 位 ABI 下（Volta / Turing / Ampere / Hopper）：

偏移地址	符号含义	读取形式示例	详细解析
`c[0x0][0x0]`	`gridDim.x`	`ULDC UR5, c[0x0][0x0]`	Grid X 维度大小（32 位无符号整数）
`c[0x0][0x4]`	`gridDim.y`	`ULDC UR6, c[0x0][0x4]`	Grid Y 维度大小
`c[0x0][0x8]`	`gridDim.z`	`ULDC UR7, c[0x0][0x8]`	Grid Z 维度大小
`c[0x0][0xc]`	`blockDim.x`	`ULDC UR4, c[0x0][0xc]`	Block X 维度大小
`c[0x0][0x10]`	`blockDim.y`	`ULDC UR5, c[0x0][0x10]`	Block Y 维度大小
`c[0x0][0x14]`	`blockDim.z`	`ULDC UR6, c[0x0][0x14]`	Block Z 维度大小
`c[0x0][0x28]`	Stack Pointer (SP)	`LDC R1, c[0x0][0x28]`	每个线程局部内存栈帧基准偏移指针（初始化 `R1`）
`c[0x0][0x208]`	Descriptor Base	`ULDC.64 UR4, c[0x0][0x208]`	Hopper 内存描述符基地址，全局寻址基点
`c[0x0][0x210]`	Kernel Param 0	`ULDC UR7, c[0x0][0x210]`	Kernel 签名的第 1 个实参（按类型大小对齐）
`c[0x0][0x218]`	Kernel Param 1	`ULDC.64 UR6, c[0x0][0x218]`	Kernel 签名的第 2 个实参（如 64 位指针）
`c[0x0][0x220+]`	Kernel Param 2+	`LDC.64 R2, c[0x0][0x220]`	后续入参，若被向量运算消费可能直接用 `LDC` 读入 `R`

4. 操作数修饰符与特殊寻址形式

.reuse 后缀（操作数缓存复用）：例如 IMAD R4, R17.reuse, UR4, RZ;。硬件在 SM 发射单元配有小容量 Operand Reuse Cache。当一条指令的输入操作数在极短周期后会被另一条指令再次作为源操作数使用时，编译期标记 .reuse。硬件调度器直接从缓存捕获该值，不再发起向物理寄存器堆（Register File）的读事务，从而消除 Register Bank Conflict 并显著节约读功耗。
desc[URx][Ry.64]（描述符基址寻址）： Hopper（SM 90+）针对全局/共享内存引入的寻址格式。
- URx：64 位的 Uniform 描述符/基地址（整 Warp 32 线程共享一份，由 ULDC.64 载入）。
- Ry.64：64 位或 32 位扩展的 Per-thread 偏移量（每个线程私有）。
- 硬件收益：传统架构下每个线程都必须在自己的向量寄存器中持有一份 64 位基地址指针（消耗 2 个 VGPR）；Hopper 通过 desc[URx][Ry.64] 将基地址抽离进标量 UR，直接为每个线程释放 2 个 VGPR。

三、成对指令对照表（核心架构速查）

NVIDIA 从 Turing（SM 7.5）起为几乎所有基础整数、逻辑、位移与比较指令引入了对应的 U 前缀版本。下表逐对对照向量通路指令与 Uniform 通路指令。

确认状态说明：表中 11 对指令均经官方《CUDA Binary Utilities》指令集规范（Turing Table 6、Ampere Table 7、Hopper Table 8）及 CUDA 12.8 真实编译验证确认存在。

向量版指令	Uniform 版指令	核心功能区别	真实 SASS 示例	对应 PTX 指令	Uniform 版生成触发条件
`S2R`	`S2UR`	读特殊寄存器到向量 `R` vs 标量 `UR`	`S2R R0, SR_TID.X` `S2UR UR4, SR_CTAID.X`	`mov.u32 %r, %ctaid.x`	特殊寄存器属于 Warp-uniform（如 `blockIdx`、`gridDim`、Hopper `SR_CgaCtaId`）
`IMAD`	`UIMAD`	32 位乘加： $d = a \times b + c$	`IMAD R4, R17, UR4, RZ` `UIMAD UR4, UR4, 0xa, URZ`	`mad.lo.s32 %r, ...`	乘数、被乘数、累加项全部来源于常量、`blockIdx` 或 uniform 循环变量
`IADD3`	`UIADD3`	3 输入 32 位整数加： $d = a + b + c$	`IADD3 R0, R0, 0x3, RZ` `UIADD3 UR4, UR4, 0x1, URZ`	连续两条加法合并（PTX 无三输入加）	三个操作数在 Warp 间全一致（如 uniform 循环步进、地址基准偏移累加）
`LDC`	`ULDC`	从 Constant Bank 载入到 `R` vs `UR`	`LDC R17, c[0x0][0x220]` `ULDC.64 UR4, c[0x0][0x208]`	`ld.param.* %rd, [...]`	参数读取后仅参与 uniform 表达式，或作为 Hopper 访存描述符基地址
`MOV`	`UMOV`	寄存器赋值 / 立即数装载	`MOV R5, UR5` `UMOV UR4, URZ`	`mov.b32 %r, ...`	将常量 0 或 uniform 变量初始化到 `UR`
`LEA`	`ULEA`	加载有效地址（带位移加法）	`LEA R0, P0, R1, R2, 0x2` `ULEA UR4, UP0, UR5, UR6, 0x2`	编译器地址折叠合成	纯标量结构体偏移或数组基地址计算
`LOP3`	`ULOP3`	任意 3 输入布尔逻辑运算（LUT）	`LOP3.LUT R5, R5, 0xe0, R0, 0xe2, !PT` `ULOP3.LUT UR4, UR4, UR5, UR6, 0xca, !UPT`	逻辑门折叠（PTX 无 3 输入逻辑）	位掩码提取、标志位合并且所有输入均为 uniform
`SHF`	`USHF`	双字漏斗移位（Funnel Shift）	`SHF.L.U32 R17, R17, 0x4, RZ` `USHF.R.S32 UR4, URZ, 0x2, UR5`	`shf.l.` / `shf.r.`	跨 64 位大整数移位或高精度定点数截断，数据流源头为 uniform
`SEL`	`USEL`	基于谓词条件二选一数据	`SEL R0, R1, R2, P0` `USEL UR4, UR5, UR6, UP0`	`selp.b32 %r, %a, %b, %p`	三元表达式 `(u_cond) ? u_a : u_b`，且条件来自 `UISETP` 产生的 `UP`
`ISETP`	`UISETP`	整数比较，写向量 `P` vs 标量 `UP`	`ISETP.NE.AND P0, PT, R0, RZ, PT` `UISETP.LT.AND UP0, UPT, UR4, 0xa, UPT`	`setp.lt.s32 %p, ...`	比较双方均为 uniform 寄存器，且结果用于控制 Warp 统一分支或 `USEL`
`PRMT`	`UPRMT`	4 字节任意排列重组（Byte Permute）	`PRMT R7, R0, 0x123, R5` `UPRMT UR4, UR5, 0x123, UR6`	PTX `prmt.b32` / `__byte_perm`	字节序翻转（Endian swap）或短整型重打包，输入数据为 uniform

四、不成对的重要指令

除成对出现的指令外，SASS 还有一批在控制流、跨 Lane 协作与跨流水线调度中起中枢作用的非对称指令：

1. 算术与地址扩展类

IMAD.WIDE / IMAD.WIDE.U32
- 硬件语义：32×32 位带符号/无符号乘法，并将结果累加到 64 位寄存器对中： $d.64 = a.32 \times b.32 + c.64$ 。
- 真实示例：IMAD.WIDE R2, R5, 0x4, R2 ;
- 为什么是不成对的核心：GPU 访存全部为 64 位虚拟地址，但线程索引与步长通常是 32 位整型。IMAD.WIDE 是一切指针算术（Pointer Math：base_ptr + idx * sizeof(T)）的绝对主力，单指令完成“乘类型大小并加到 64 位基址”。
- 注意：Hopper 中也存在 uniform 版 UIMAD.WIDE UR4, UR6, 0x4, UR4，用于计算整个 Warp 共享的张量切片 64 位基地址。

2. Warp 级数据交换与规约

SHFL.IDX / SHFL.UP / SHFL.DOWN / SHFL.BFLY
- 硬件语义：Warp 内跨 Lane 寄存器无显存直换。
  - .IDX：任意 Lane 寻址广播（对应 __shfl_sync）。
  - .UP：向高编号 Lane 传递（laneId - delta，对应 __shfl_up_sync）。
  - .DOWN：向低编号 Lane 传递（laneId + delta，对应 __shfl_down_sync）。
  - .BFLY：蝶式异或交换（laneId ^ mask，对应 __shfl_xor_sync）。
- 真实示例：SHFL.DOWN PT, R3, R7, 0x1, 0x1f ;（从 laneId + 1 搬移到当前 Lane，mask 为全 32 线程）。
REDUX（硬件 Warp 归约，SM 80+）
- 硬件语义：将 Warp 内 32 个线程的向量寄存器直接进行加/最小值/按位规约，并将最终结果直接写入标量 UR！
- 真实示例：REDUX.SUM.S32 UR4, R7 ;
- 架构演进：SM 7.5 及之前需要 5 轮 SHFL.DOWN + IADD 软件拼接；Ampere 起由单条 REDUX 硬件指令一口气吞下，且结果天然落在标量通路，无缝衔接后续控制流。
ELECT（硬件选主，SM 90+）
- 硬件语义：在当前活跃线程集合中，以硬件仲裁选出唯一一个 Leader 线程（通常为当前活跃的最低编号 Lane）。
- 真实示例：ELECT P0, URZ, PT ;
- 作用：被选中的线程其 P0 谓词置 True，其余活跃线程置 False。在 Hopper TMA（Tensor Memory Accelerator）发起的流水线中，用于单线程发起异步搬运，彻底省去以往 if (laneId == 0) 的算术比较开销。

3. Uniform 专用统计与位操作

VOTEU.ANY / VOTEU.ALL
- 硬件语义：对 Warp 内部所有线程的条件谓词进行投票聚合，结果直接写回 UR 或 UP。
- 真实示例：VOTEU.ANY UR4, UPT, PT ;
- 对比：常规 VOTE 写入每线程的 R 或 P；VOTEU 将汇总结果直送 Uniform 通路，供后续标量跳转决策。
UFLO.U32
- 硬件语义：Uniform Find-Leading-One。在 32 位 UR 中检索最高有效位（MSB）位置。
- 用途：标量除法定点近似、动态共享内存地址对齐与规格化。
UPOPC
- 硬件语义：Uniform Population Count。计算 UR 中二进制位为 1 的总个数。

4. 向量与 Uniform 寄存器间搬运：`R2UR` 与 `MOV R, UR`

R2UR（Vector Register → Uniform Register）
- 官方定义：Move from Vector Register to a Uniform Register。
- 硬件机制：从 32 个向量 Lane 中采样（通常读取 Lane 0）或按规则折叠收敛写入单个 UR。
从 UR 搬运回 R 的官方机制：没有名为 UR2R 的独立 Opcode！
- 查证确认：查阅 NVIDIA 官方 Turing/Ampere/Hopper 指令集表，不存在名为 UR2R 的汇编助记符。
- 实际执行：在硬件上，所有常规向量指令天然支持将 UR 作为源操作数。若需单纯的数据转移，编译器直接发射 MOV R, UR（如 MOV R5, UR5 ;）或 IMAD.U32 R, RZ, RZ, UR。
- 工程诊断核心：如果在反汇编中看到密集的 R2UR 与 MOV R, UR，表明程序中原本以为是标量的数据流发生了断裂，编译器无法全程在标量寄存器堆闭环，被迫在两套物理 Register File 之间反复来回搬运，产生额外的发射停顿。

5. 访存与控制流

访存体系：
- LDG.E / STG.E：64 位扩展全局内存读取与写入。
- LDG.E.CONSTANT：带常量/纹理只读缓存提示的全局内存读取。
- LDS / STS：共享内存（Shared Memory）读写。
- LDC / ULDC：从 Constant Bank 读取数据。
控制流体系：
- BSSY B0, label：设置独立线程调度汇聚屏障栈（Barrier Set Synchronization）。
- BSYNC B0：在分支交汇处同步屏障上的所有线程（Barrier Synchronization）。
- BRA label：相对分支跳转（可受谓词控制 @P0 BRA 或 @UP0 BRA）。
- EXIT：当前线程/CTA 终止执行并释放资源。
- NOP：空操作（常用于填充发射间隙或对齐指令字节边界）。

五、Warp 级 PTX 原语速查

下表列出 CUDA 算子开发中最常用的 Warp 级 PTX 汇编原语、对应 C++ 内建接口、最低计算能力与编译下沉目标：

PTX 原语语法	CUDA C++ 内建函数 / 接口	最低架构	对应 SASS 汇编	典型用途	注意事项与陷阱
`shfl.sync.aligned.b32 d\|p, a, b, c, mask;`	`__shfl_sync(mask, val, srcLane)` `__shfl_down_sync` `__shfl_up_sync` `__shfl_xor_sync`	sm_70 (注1)	`SHFL.IDX` `SHFL.DOWN` `SHFL.UP` `SHFL.BFLY`	Warp 内规约、滑窗卷积、广播通信	掩码 `mask` 内的线程必须同时到达，非 mask 线程读取结果未定义
`vote.sync.aligned.pred d, p, mask;` `vote.sync.ballot.b32 d, p, mask;`	`__all_sync(mask, pred)` `__any_sync(mask, pred)` `__ballot_sync(mask, pred)`	sm_70 (注1)	`VOTE.ALL` `VOTE.ANY` `VOTE` / `VOTEU`	循环收敛判定、条件稀疏跳过、位掩码生成	`ballot` 返回 32 位整型，每一位对应一个 Lane 的谓词真假
`redux.sync.aligned.op.type d, a, mask;`	`__reduce_add_sync(mask, val)` `__reduce_min_sync` `__reduce_max_sync` 等	sm_80	`REDUX.SUM` `REDUX.MIN` `REDUX.MAX`	Warp 全局聚合，无需多轮 Shuffle 展开	必须全 Warp 32 线程参与（`mask = 0xffffffff`）；仅限整型运算
`elect.sync d\|p, mask;`	PTX 内联汇编 / CUTLASS 3.x `cute::elect_one_sync()`	sm_90	`ELECT P0, URZ, PT`	TMA 异步传输单线程选主、Barrier 触发	仅由硬件挑选一个 Leader，结果为 1 位谓词，不带来任何分支同步开销
`activemask.b32 d;`	`__activemask()`	sm_30	特殊寄存器读取或编译器活跃追踪	运行时动态查询当前到达该指令的时钟活跃 Lane 掩码	仅表示瞬时活跃，不能保证后续指令的收敛性，不能取代显式 `mask`
`bar.warp.sync mask;`	`__syncwarp(mask = 0xffffffff)`	sm_70	`WARPSYNC` / `BSYNC`	强制 Warp 内指定线程在此硬同步，形成执行屏障	如果 `mask` 包含了在其他分支永远无法到达该点的线程，将引发死锁
`match.any.sync.b32 d, a, mask;` `match.all.sync.b32 d\|p, a, mask;`	`__match_any_sync(mask, val)` `__match_all_sync(mask, val)`	sm_70	`MATCH.ANY` `MATCH.ALL`	跨线程值匹配、同值聚合（Value Clustering）、基数排序	快速找出 Warp 内部拥有相同数据的线程集合掩码
`bra.uni target;` `setp.uni.cmp.type p, a, b;`	编译器自动分析或 PTX 内联带 `.uni`	sm_70	`@UP0 BRA` / `@P0 BRA`	无分歧的统一跳转与统一条件判断	`.uni` 显式保证所有线程跳转方向 100% 一致，通知 ptxas 免除设置 `BSSY` 屏障栈

(注1)：无 .sync 的旧版 shfl 与 vote 在早期架构（如 sm_30/sm_12）即可用，但自 Volta（sm_70）引入 Independent Thread Scheduling 机制后，无掩码版本全部弃用，强制要求 .sync 保证收敛。

六、实战：怎么读一段真实 SASS

下面这段汇编截取自我们在 Hopper（sm_90）集群上编译 t2.cu 中 kB_cta 内核的真实输出（对应二维行偏移张量访存）：

/*0010*/ S2UR UR4, SR_CTAID.Y ;                        // 1. 将 blockIdx.y 读取到标量寄存器 UR4
/*0020*/ S2R  R0,  SR_TID.X ;                          // 2. 将 threadIdx.x 读取到向量寄存器 R0
/*0050*/ IMAD R4, R17.reuse, UR4, RZ ;                 // 3. 向量乘加：R4 = R17(row_stride) * UR4 + 0
/*0060*/ SHF.L.U32 R17, R17, 0x4, RZ ;                 // 4. 向量逻辑左移：R17 = R17 << 4 (乘以 16)
/*0070*/ ULDC.64 UR4, c[0x0][0x208] ;                  // 5. 标量载入：从常数内存 0x208 装载 64 位描述符基址到 UR4
/*0080*/ SHF.L.U32 R5, R4, 0x5, RZ ;                   // 6. 向量逻辑左移：R5 = R4 << 5 (乘以 32)
/*0090*/ LOP3.LUT R5, R5, 0xffffffe0, R0, 0xe2, !PT ;  // 7. 三输入位运算：拼接基址与 (threadIdx.x & 31)
/*00a0*/ IMAD.WIDE R2, R5, 0x4, R2 ;                   // 8. 拓宽乘加：R2.64 = R5 * 4 字节 + R2(in 基准偏移)
/*00c0*/ LDG.E.CONSTANT R2, desc[UR4][R2.64] ;         // 9. 描述符全局只读访存：从 desc[UR4] + R2.64 读取 float

深度逐行解码

S2UR UR4, SR_CTAID.Y
- 编译器静态推断 blockIdx.y 是整个 Warp 严格相等的量，直接走标量通路 S2UR 放入 UR4，不消耗任何 per-thread VGPR。
S2R R0, SR_TID.X
- threadIdx.x 各 Lane 互不相同（Divergent），走向量通路 S2R 分配到 R0。
IMAD R4, R17.reuse, UR4, RZ（核心洞察：向量指令直接读 UR）
- 指令性质：这是一条向量 IMAD 指令（目标是向量寄存器 R4），并非 UIMAD！
- 关键认知：数据放入 UR 并不等于后续计算走 Uniform Datapath！
- 这里的 UR4 作为向量乘法的一个源操作数，硬件在发射时自动将 UR4 的单份标量值广播到 SIMT 流水线的 32 个执行通道。R17.reuse 表示复用操作数缓存，乘完后存入每个线程私有的 R4。
SHF.L.U32 R17, R17, 0x4, RZ
- 预计算下一循环迭代的跨步增量（ $16 \times \text{row\_stride}$ ），通过左移 4 位实现。
ULDC.64 UR4, c[0x0][0x208]
- 在访存前，将全局张量描述符基地址装入 64 位的 UR4（覆盖了之前已使用完的 blockIdx.y，发生寄存器复用）。
SHF.L.U32 R5, R4, 0x5, RZ
- 将 R4 左移 5 位（即乘以 32），计算基准行偏移的元素数量。
LOP3.LUT R5, R5, 0xffffffe0, R0, 0xe2, !PT
- 经典比特折叠操作。通过查找表掩码 0xe2，等价于位运算表达式 (R5 & 0xffffffe0) | (R0 & 0x1f)，在一条指令内完成了将“行基准偏移”与每个线程私有的 lane = threadIdx.x & 31 完美合成！
IMAD.WIDE R2, R5, 0x4, R2
- 将元素索引 R5 乘以 4（sizeof(float)），并加上参数传入的基准指针偏移 R2，拓宽为 64 位字节偏移量存入 R2–R3 向量寄存器对。
LDG.E.CONSTANT R2, desc[UR4][R2.64]
- Hopper 现代访存语法：硬件将 Warp 统一持有的 desc[UR4] 描述符基地址与每线程计算出的 R2.64 字节偏移在 AGU（地址生成单元）中直接加和，从只读常量缓存发起并发内存事务。

七、SASS 读码诊断清单（Diagnostic Checklist）

当你在性能分析器（NCU）或反汇编中审计 SASS 时，请重点关注以下四类关键信号：

信号 1：出现大量 `R2UR` 或 `MOV R, UR` 跨通路搬运

表象：代码中密集插入了 R2UR UR_x, R_y 或 MOV R_x, UR_y。
诊断意义：手工暴露 Uniformity 可能完全白干甚至起反作用！
机理：这说明你在 CUDA 代码中虽然将某些计算拆成了“看起来 uniform”的形式，但数据流的前驱或后继依赖依然锚定在向量通路上。编译器无法维持端到端的 Uniform 流水线，被迫在向量 Register File 与标量 Register File 之间频繁插装数据搬运，带来额外的指令发射开销与跨寄存器堆延迟。

信号 2：`UIMAD` / `UIADD3` 为 0，但指令中大量出现 `UR` 操作数

表象：汇编中没有任何 U 前缀的算术指令，但常规向量指令（如 IMAD、FADD、IADD3）的源操作数里全是 UR。
诊断意义：值进了标量寄存器，但计算依然全堆在向量流水线上。
机理：UR 在这里仅被编译器当作了“常量广播源”（节省了 VGPR 存储）。但如果你优化的初衷是“利用 Uniform Datapath 卸载向量 ALU 负载、提高算术吞吐”，这一目标并未实现。需要检查下游消费该变量的操作数是否存在发散项。

信号 3：本该 Uniform 的变量使用了 `S2R` 而非 `S2UR`

表象：读取 blockIdx、gridDim 或全 Warp 统一的入参时，SASS 发射了 S2R R_x, SR_CTAID.X 而不是 S2UR。
诊断意义：编译器 Divergence Analysis 失败，未能证明 Uniformity。
根因排查：倒查该变量的数据流祖先节点。在 CUDA C++ 模板展开或内联函数中，是否在某些分支路径混入了与 threadIdx 相关的运算，或者指针别名（Aliasing）导致编译器假定内存可能被某一特定线程修改，从而保守降级为向量处理。

信号 4：`-Xptxas -v` 报告的寄存器数跨越关键门槛

表象：nvcc -Xptxas -v 显示单个线程使用的寄存器数从 128 变成 129，或从 168 变成 170。
诊断意义：Occupancy 触发阶梯式断崖崩溃！
硬件分界点（以现代 SM 架构 64K VGPR 为准）：
- $\le 32$ 寄存器：满 Occupancy（理论可驻留最大 Warp 数，如 48/64 Warps）；
- $33 \sim 64$ 寄存器：中高 Occupancy；
- $65 \sim 128$ 寄存器：最大只能跑 16 个 Warp（50% Occupancy）；
- 跨过 128（如 129）：单 SM 驻留 Warp 数断崖跌至 8 个（25% Occupancy）；
- 跨过 255：触发溢出，进入 Local Memory（显存/L1 Cache Spill），访存延迟暴增。
对策：审计反汇编中的 64 位基地址指针与循环累加变量，检查能否通过调整算式顺序，让 ptxas 能够将其折叠到 UR0–UR63 中，将 VGPR 压回关键门槛之内。

八、附录：确认状态与参考出处

1. 指令存在性确认汇总

汇编指令	官方 ISA 规范状态	出处与验证依据	架构支持范围
`S2UR` / `S2R`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`UIMAD` / `IMAD`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`UIADD3` / `IADD3`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`ULDC` / `LDC`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`UMOV` / `MOV`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`ULEA` / `LEA`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`ULOP3` / `LOP3`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`USHF` / `SHF`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`USEL` / `SEL`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`UISETP` / `ISETP`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`UPRMT` / `PRMT`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`R2UR`	完全确认	NVIDIA CUDA Binary Utilities (Table 6, 7, 8)	Turing (sm_75) 及以上
`UR2R`	未确认（无独立指令名）	官方表中无此指令；硬件由常规向量指令（如 `MOV R, UR`）直接读取	概念称谓，非独立 Opcode
`REDUX`	完全确认	NVIDIA CUDA Binary Utilities (Table 7, 8)	Ampere (sm_80) 及以上
`ELECT`	完全确认	NVIDIA CUDA Binary Utilities (Table 8)	Hopper (sm_90) 及以上

2. 官方参考规范

CUDA Binary Utilities Reference Guide (NVIDIA Documentation): https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html
- Section 4.1: Turing Instruction Set (Table 6)
- Section 4.2: Ampere and Ada Instruction Set (Table 7)
- Section 4.3: Hopper Instruction Set (Table 8)
Parallel Thread Execution ISA Version 8.7 / 8.8 (NVIDIA Documentation): https://docs.nvidia.com/cuda/parallel-thread-execution/index.html
- Warp Shuffle Functions (shfl.sync)
- Warp Vote Functions (vote.sync)
- Warp Reduction Functions (redux.sync)
- Leader Election Functions (elect.sync)
CUDA C++ Programming Guide: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
- Section B.15: Warp Shuffle Functions
- Section B.16: Warp Vote Functions
- Section B.17: Warp Match Functions
- Section B.18: Warp Reduce Functions