6.4 组 4：访存与原子（unit: lsu，301 条）
cache 维护（3 条）
CACHE
语法（Syntax）
cache.inv.l1
cache.inv.l2
cache.wb.l2
| 变体 | 编码 | 延迟 | 语义 | 标志 |
|---|---|---|---|---|
cache.inv.l1 | SYS op opcode=0x4e f5=0x9 | async:vm | 无效化本 half-core 的 L1D 与纹理 cache（device 级 acquire） | ei co |
cache.inv.l2 | SYS op opcode=0x4e f5=0xa | async:vm | 无效化 L2 中非一致 host 内存行（system 级 acquire） | ei co |
cache.wb.l2 | SYS op opcode=0x4e f5=0xb | async:vst | 回写非一致 host 内存的脏 L2 行（system 级 release） | ei co |
向量加载（40 条）
LD
语法（Syntax）
ld.global.u8 rd:v rs1:vp^2 imm:s24 cpol:cpol
ld.global.i8 rd:v rs1:vp^2 imm:s24 cpol:cpol
ld.global.u16 rd:v rs1:vp^2 imm:s24 cpol:cpol
ld.global.i16 rd:v rs1:vp^2 imm:s24 cpol:cpol
ld.global.b32 rd:v rs1:vp^2 imm:s24 cpol:cpol
ld.global.b64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
ld.global.b96 rd:v^3 rs1:vp^2 imm:s24 cpol:cpol
ld.global.b128 rd:v^4 rs1:vp^2 imm:s24 cpol:cpol
ld.ubase.u8 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.ubase.i8 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.ubase.u16 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.ubase.i16 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.ubase.b32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.ubase.b64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.ubase.b96 rd:v^3 rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.ubase.b128 rd:v^4 rs2:up^2 rs1:v imm:s24 cpol:cpol
ld.buf.u8 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.buf.i8 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.buf.u16 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.buf.i16 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.buf.b32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.buf.b64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.buf.b96 rd:v^3 rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.buf.b128 rd:v^4 rs2:d4^4 rs1:v imm:s24 cpol:cpol
ld.shared.u8 rd:v rs1:v imm:s24
ld.shared.i8 rd:v rs1:v imm:s24
ld.shared.u16 rd:v rs1:v imm:s24
ld.shared.i16 rd:v rs1:v imm:s24
ld.shared.b32 rd:v rs1:v imm:s24
ld.shared.b64 rd:v^2 rs1:v imm:s24
ld.shared.b96 rd:v^3 rs1:v imm:s24
ld.shared.b128 rd:v^4 rs1:v imm:s24
ld.scratch.u8 rd:v rs1:v imm:s24
ld.scratch.i8 rd:v rs1:v imm:s24
ld.scratch.u16 rd:v rs1:v imm:s24
ld.scratch.i16 rd:v rs1:v imm:s24
ld.scratch.b32 rd:v rs1:v imm:s24
ld.scratch.b64 rd:v^2 rs1:v imm:s24
ld.scratch.b96 rd:v^3 rs1:v imm:s24
ld.scratch.b128 rd:v^4 rs1:v imm:s24
| 变体 | 编码 | 延迟 | 语义 | 标志 |
|---|---|---|---|---|
ld.global.u8 | M op opcode=0x15 f5=0x0 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 零扩展读取 8 位 | ld |
ld.global.i8 | M op opcode=0x15 f5=0x1 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 符号扩展读取 8 位 | ld |
ld.global.u16 | M op opcode=0x15 f5=0x2 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 零扩展读取 16 位 | ld |
ld.global.i16 | M op opcode=0x15 f5=0x3 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 符号扩展读取 16 位 | ld |
ld.global.b32 | M op opcode=0x15 f5=0x4 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 读取 32 位 | ld |
ld.global.b64 | M op opcode=0x15 f5=0x5 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 读取 64 位 | ld |
ld.global.b96 | M op opcode=0x15 f5=0x6 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 读取 96 位 | ld |
ld.global.b128 | M op opcode=0x15 f5=0x7 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；d = 读取 128 位 | ld |
ld.ubase.u8 | M op opcode=0x15 f5=0x0 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 零扩展读取 8 位 | ld |
ld.ubase.i8 | M op opcode=0x15 f5=0x1 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 符号扩展读取 8 位 | ld |
ld.ubase.u16 | M op opcode=0x15 f5=0x2 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 零扩展读取 16 位 | ld |
ld.ubase.i16 | M op opcode=0x15 f5=0x3 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 符号扩展读取 16 位 | ld |
ld.ubase.b32 | M op opcode=0x15 f5=0x4 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 读取 32 位 | ld |
ld.ubase.b64 | M op opcode=0x15 f5=0x5 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 读取 64 位 | ld |
ld.ubase.b96 | M op opcode=0x15 f5=0x6 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 读取 96 位 | ld |
ld.ubase.b128 | M op opcode=0x15 f5=0x7 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；d = 读取 128 位 | ld |
ld.buf.u8 | M op opcode=0x15 f5=0x0 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 零扩展读取 8 位；越界 / null 读 0 | ld bc |
ld.buf.i8 | M op opcode=0x15 f5=0x1 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 符号扩展读取 8 位；越界 / null 读 0 | ld bc |
ld.buf.u16 | M op opcode=0x15 f5=0x2 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 零扩展读取 16 位；越界 / null 读 0 | ld bc |
ld.buf.i16 | M op opcode=0x15 f5=0x3 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 符号扩展读取 16 位；越界 / null 读 0 | ld bc |
ld.buf.b32 | M op opcode=0x15 f5=0x4 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 读取 32 位；越界 / null 读 0 | ld bc |
ld.buf.b64 | M op opcode=0x15 f5=0x5 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 读取 64 位；越界 / null 读 0 | ld bc |
ld.buf.b96 | M op opcode=0x15 f5=0x6 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 读取 96 位；越界 / null 读 0 | ld bc |
ld.buf.b128 | M op opcode=0x15 f5=0x7 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；d = 读取 128 位；越界 / null 读 0 | ld bc |
ld.shared.u8 | M op opcode=0x15 f5=0x0 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 零扩展读取 8 位 | ld |
ld.shared.i8 | M op opcode=0x15 f5=0x1 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 符号扩展读取 8 位 | ld |
ld.shared.u16 | M op opcode=0x15 f5=0x2 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 零扩展读取 16 位 | ld |
ld.shared.i16 | M op opcode=0x15 f5=0x3 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 符号扩展读取 16 位 | ld |
ld.shared.b32 | M op opcode=0x15 f5=0x4 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 读取 32 位 | ld |
ld.shared.b64 | M op opcode=0x15 f5=0x5 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 读取 64 位 | ld |
ld.shared.b96 | M op opcode=0x15 f5=0x6 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 读取 96 位 | ld |
ld.shared.b128 | M op opcode=0x15 f5=0x7 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；d = 读取 128 位 | ld |
ld.scratch.u8 | M op opcode=0x15 f5=0x0 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 零扩展读取 8 位 | ld |
ld.scratch.i8 | M op opcode=0x15 f5=0x1 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 符号扩展读取 8 位 | ld |
ld.scratch.u16 | M op opcode=0x15 f5=0x2 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 零扩展读取 16 位 | ld |
ld.scratch.i16 | M op opcode=0x15 f5=0x3 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 符号扩展读取 16 位 | ld |
ld.scratch.b32 | M op opcode=0x15 f5=0x4 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 读取 32 位 | ld |
ld.scratch.b64 | M op opcode=0x15 f5=0x5 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 读取 64 位 | ld |
ld.scratch.b96 | M op opcode=0x15 f5=0x6 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 读取 96 位 | ld |
ld.scratch.b128 | M op opcode=0x15 f5=0x7 sp=0x4 f2=0x0 | async:vm | A = lane 私有 scratch 地址 (s1 + sext(imm))；d = 读取 128 位 | ld |
向量存储（30 条）
ST
语法（Syntax）
st.global.b8 rd:v rs1:vp^2 imm:s24 cpol:cpol
st.global.b16 rd:v rs1:vp^2 imm:s24 cpol:cpol
st.global.b32 rd:v rs1:vp^2 imm:s24 cpol:cpol
st.global.b64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
st.global.b96 rd:v^3 rs1:vp^2 imm:s24 cpol:cpol
st.global.b128 rd:v^4 rs1:vp^2 imm:s24 cpol:cpol
st.ubase.b8 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
st.ubase.b16 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
st.ubase.b32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
st.ubase.b64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
st.ubase.b96 rd:v^3 rs2:up^2 rs1:v imm:s24 cpol:cpol
st.ubase.b128 rd:v^4 rs2:up^2 rs1:v imm:s24 cpol:cpol
st.buf.b8 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
st.buf.b16 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
st.buf.b32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
st.buf.b64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
st.buf.b96 rd:v^3 rs2:d4^4 rs1:v imm:s24 cpol:cpol
st.buf.b128 rd:v^4 rs2:d4^4 rs1:v imm:s24 cpol:cpol
st.shared.b8 rd:v rs1:v imm:s24
st.shared.b16 rd:v rs1:v imm:s24
st.shared.b32 rd:v rs1:v imm:s24
st.shared.b64 rd:v^2 rs1:v imm:s24
st.shared.b96 rd:v^3 rs1:v imm:s24
st.shared.b128 rd:v^4 rs1:v imm:s24
st.scratch.b8 rd:v rs1:v imm:s24
st.scratch.b16 rd:v rs1:v imm:s24
st.scratch.b32 rd:v rs1:v imm:s24
st.scratch.b64 rd:v^2 rs1:v imm:s24
st.scratch.b96 rd:v^3 rs1:v imm:s24
st.scratch.b128 rd:v^4 rs1:v imm:s24
| 变体 | 编码 | 延迟 | 语义 | 标志 |
|---|---|---|---|---|
st.global.b8 | M op opcode=0x19 f5=0x0 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；将数据寄存器低 8 位写入 A | st |
st.global.b16 | M op opcode=0x19 f5=0x1 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；将数据寄存器低 16 位写入 A | st |
st.global.b32 | M op opcode=0x19 f5=0x2 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；将数据寄存器低 32 位写入 A | st |
st.global.b64 | M op opcode=0x19 f5=0x3 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；将数据寄存器低 64 位写入 A | st |
st.global.b96 | M op opcode=0x19 f5=0x4 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；将数据寄存器低 96 位写入 A | st |
st.global.b128 | M op opcode=0x19 f5=0x5 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；将数据寄存器低 128 位写入 A | st |
st.ubase.b8 | M op opcode=0x19 f5=0x0 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；将数据寄存器低 8 位写入 A | st |
st.ubase.b16 | M op opcode=0x19 f5=0x1 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；将数据寄存器低 16 位写入 A | st |
st.ubase.b32 | M op opcode=0x19 f5=0x2 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；将数据寄存器低 32 位写入 A | st |
st.ubase.b64 | M op opcode=0x19 f5=0x3 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；将数据寄存器低 64 位写入 A | st |
st.ubase.b96 | M op opcode=0x19 f5=0x4 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；将数据寄存器低 96 位写入 A | st |
st.ubase.b128 | M op opcode=0x19 f5=0x5 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；将数据寄存器低 128 位写入 A | st |
st.buf.b8 | M op opcode=0x19 f5=0x0 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；将数据寄存器低 8 位写入 A | st bc |
st.buf.b16 | M op opcode=0x19 f5=0x1 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；将数据寄存器低 16 位写入 A | st bc |
st.buf.b32 | M op opcode=0x19 f5=0x2 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；将数据寄存器低 32 位写入 A | st bc |
st.buf.b64 | M op opcode=0x19 f5=0x3 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；将数据寄存器低 64 位写入 A | st bc |
st.buf.b96 | M op opcode=0x19 f5=0x4 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；将数据寄存器低 96 位写入 A | st bc |
st.buf.b128 | M op opcode=0x19 f5=0x5 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；将数据寄存器低 128 位写入 A | st bc |
st.shared.b8 | M op opcode=0x19 f5=0x0 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；将数据寄存器低 8 位写入 A | st |
st.shared.b16 | M op opcode=0x19 f5=0x1 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；将数据寄存器低 16 位写入 A | st |
st.shared.b32 | M op opcode=0x19 f5=0x2 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；将数据寄存器低 32 位写入 A | st |
st.shared.b64 | M op opcode=0x19 f5=0x3 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；将数据寄存器低 64 位写入 A | st |
st.shared.b96 | M op opcode=0x19 f5=0x4 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；将数据寄存器低 96 位写入 A | st |
st.shared.b128 | M op opcode=0x19 f5=0x5 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；将数据寄存器低 128 位写入 A | st |
st.scratch.b8 | M op opcode=0x19 f5=0x0 sp=0x4 f2=0x0 | async:vst | A = lane 私有 scratch 地址 (s1 + sext(imm))；将数据寄存器低 8 位写入 A | st |
st.scratch.b16 | M op opcode=0x19 f5=0x1 sp=0x4 f2=0x0 | async:vst | A = lane 私有 scratch 地址 (s1 + sext(imm))；将数据寄存器低 16 位写入 A | st |
st.scratch.b32 | M op opcode=0x19 f5=0x2 sp=0x4 f2=0x0 | async:vst | A = lane 私有 scratch 地址 (s1 + sext(imm))；将数据寄存器低 32 位写入 A | st |
st.scratch.b64 | M op opcode=0x19 f5=0x3 sp=0x4 f2=0x0 | async:vst | A = lane 私有 scratch 地址 (s1 + sext(imm))；将数据寄存器低 64 位写入 A | st |
st.scratch.b96 | M op opcode=0x19 f5=0x4 sp=0x4 f2=0x0 | async:vst | A = lane 私有 scratch 地址 (s1 + sext(imm))；将数据寄存器低 96 位写入 A | st |
st.scratch.b128 | M op opcode=0x19 f5=0x5 sp=0x4 f2=0x0 | async:vst | A = lane 私有 scratch 地址 (s1 + sext(imm))；将数据寄存器低 128 位写入 A | st |
原子（无返回）（104 条）
ATOM
语法（Syntax）
atom.global.add.u32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.sub.u32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.and.b32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.or.b32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.xor.b32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.i32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.i32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.u32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.u32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.xchg.b32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.cmpxchg.b32 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.inc.u32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.dec.u32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.fadd.f32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.fmin.f32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.fmax.f32 rd:v rs1:vp^2 imm:s24 cpol:cpol
atom.global.add.u64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.and.b64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.or.b64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.xor.b64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.i64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.i64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.u64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.u64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.xchg.b64 rd:v^2 rs1:vp^2 imm:s24 cpol:cpol
atom.global.cmpxchg.b64 rd:v^4 rs1:vp^2 imm:s24 cpol:cpol
atom.ubase.add.u32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.sub.u32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.and.b32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.or.b32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xor.b32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.i32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.i32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.u32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.u32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xchg.b32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.cmpxchg.b32 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.inc.u32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.dec.u32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.fadd.f32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.fmin.f32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.fmax.f32 rd:v rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.add.u64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.and.b64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.or.b64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xor.b64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.i64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.i64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.u64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.u64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xchg.b64 rd:v^2 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.cmpxchg.b64 rd:v^4 rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.buf.add.u32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.sub.u32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.and.b32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.or.b32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xor.b32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.i32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.i32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.u32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.u32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xchg.b32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.cmpxchg.b32 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.inc.u32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.dec.u32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.fadd.f32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.fmin.f32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.fmax.f32 rd:v rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.add.u64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.and.b64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.or.b64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xor.b64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.i64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.i64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.u64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.u64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xchg.b64 rd:v^2 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.cmpxchg.b64 rd:v^4 rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.shared.add.u32 rd:v rs1:v imm:s24
atom.shared.sub.u32 rd:v rs1:v imm:s24
atom.shared.and.b32 rd:v rs1:v imm:s24
atom.shared.or.b32 rd:v rs1:v imm:s24
atom.shared.xor.b32 rd:v rs1:v imm:s24
atom.shared.min.i32 rd:v rs1:v imm:s24
atom.shared.max.i32 rd:v rs1:v imm:s24
atom.shared.min.u32 rd:v rs1:v imm:s24
atom.shared.max.u32 rd:v rs1:v imm:s24
atom.shared.xchg.b32 rd:v rs1:v imm:s24
atom.shared.cmpxchg.b32 rd:v^2 rs1:v imm:s24
atom.shared.inc.u32 rd:v rs1:v imm:s24
atom.shared.dec.u32 rd:v rs1:v imm:s24
atom.shared.fadd.f32 rd:v rs1:v imm:s24
atom.shared.fmin.f32 rd:v rs1:v imm:s24
atom.shared.fmax.f32 rd:v rs1:v imm:s24
atom.shared.add.u64 rd:v^2 rs1:v imm:s24
atom.shared.and.b64 rd:v^2 rs1:v imm:s24
atom.shared.or.b64 rd:v^2 rs1:v imm:s24
atom.shared.xor.b64 rd:v^2 rs1:v imm:s24
atom.shared.min.i64 rd:v^2 rs1:v imm:s24
atom.shared.max.i64 rd:v^2 rs1:v imm:s24
atom.shared.min.u64 rd:v^2 rs1:v imm:s24
atom.shared.max.u64 rd:v^2 rs1:v imm:s24
atom.shared.xchg.b64 rd:v^2 rs1:v imm:s24
atom.shared.cmpxchg.b64 rd:v^4 rs1:v imm:s24
| 变体 | 编码 | 延迟 | 语义 | 标志 |
|---|---|---|---|---|
atom.global.add.u32 | M op opcode=0x1d f5=0x0 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] + data | at |
atom.global.sub.u32 | M op opcode=0x1d f5=0x1 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] - data | at |
atom.global.and.b32 | M op opcode=0x1d f5=0x2 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] & data | at |
atom.global.or.b32 | M op opcode=0x1d f5=0x3 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] | data | at |
atom.global.xor.b32 | M op opcode=0x1d f5=0x4 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] ^ data | at |
atom.global.min.i32 | M op opcode=0x1d f5=0x5 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data) | at |
atom.global.max.i32 | M op opcode=0x1d f5=0x6 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data) | at |
atom.global.min.u32 | M op opcode=0x1d f5=0x7 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data) | at |
atom.global.max.u32 | M op opcode=0x1d f5=0x8 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data) | at |
atom.global.xchg.b32 | M op opcode=0x1d f5=0x9 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = data（交换） | at |
atom.global.cmpxchg.b32 | M op opcode=0x1d f5=0xa sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换） | at |
atom.global.inc.u32 | M op opcode=0x1d f5=0xb sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = (old >= data) ? 0 : old + 1（回绕递增） | at |
atom.global.dec.u32 | M op opcode=0x1d f5=0xc sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减） | at |
atom.global.fadd.f32 | M op opcode=0x1d f5=0xd sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] + data（binary32 原子加，固定 RNE） | at opt |
atom.global.fmin.f32 | M op opcode=0x1d f5=0xe sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = minimumNumber(M[A], data) | at opt |
atom.global.fmax.f32 | M op opcode=0x1d f5=0xf sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = maximumNumber(M[A], data) | at opt |
atom.global.add.u64 | M op opcode=0x1d f5=0x10 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] + data | at |
atom.global.and.b64 | M op opcode=0x1d f5=0x11 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] & data | at |
atom.global.or.b64 | M op opcode=0x1d f5=0x12 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] | data | at |
atom.global.xor.b64 | M op opcode=0x1d f5=0x13 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] ^ data | at |
atom.global.min.i64 | M op opcode=0x1d f5=0x14 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data) | at |
atom.global.max.i64 | M op opcode=0x1d f5=0x15 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data) | at |
atom.global.min.u64 | M op opcode=0x1d f5=0x16 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data) | at |
atom.global.max.u64 | M op opcode=0x1d f5=0x17 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data) | at |
atom.global.xchg.b64 | M op opcode=0x1d f5=0x18 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：M[A] = data（交换） | at |
atom.global.cmpxchg.b64 | M op opcode=0x1d f5=0x19 sp=0x0 f2=0x0 | async:vst | A = s1_64 + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换） | at |
atom.ubase.add.u32 | M op opcode=0x1d f5=0x0 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] + data | at |
atom.ubase.sub.u32 | M op opcode=0x1d f5=0x1 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] - data | at |
atom.ubase.and.b32 | M op opcode=0x1d f5=0x2 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] & data | at |
atom.ubase.or.b32 | M op opcode=0x1d f5=0x3 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] | data | at |
atom.ubase.xor.b32 | M op opcode=0x1d f5=0x4 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] ^ data | at |
atom.ubase.min.i32 | M op opcode=0x1d f5=0x5 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data) | at |
atom.ubase.max.i32 | M op opcode=0x1d f5=0x6 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data) | at |
atom.ubase.min.u32 | M op opcode=0x1d f5=0x7 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data) | at |
atom.ubase.max.u32 | M op opcode=0x1d f5=0x8 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data) | at |
atom.ubase.xchg.b32 | M op opcode=0x1d f5=0x9 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = data（交换） | at |
atom.ubase.cmpxchg.b32 | M op opcode=0x1d f5=0xa sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换） | at |
atom.ubase.inc.u32 | M op opcode=0x1d f5=0xb sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = (old >= data) ? 0 : old + 1（回绕递增） | at |
atom.ubase.dec.u32 | M op opcode=0x1d f5=0xc sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减） | at |
atom.ubase.fadd.f32 | M op opcode=0x1d f5=0xd sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] + data（binary32 原子加，固定 RNE） | at opt |
atom.ubase.fmin.f32 | M op opcode=0x1d f5=0xe sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = minimumNumber(M[A], data) | at opt |
atom.ubase.fmax.f32 | M op opcode=0x1d f5=0xf sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = maximumNumber(M[A], data) | at opt |
atom.ubase.add.u64 | M op opcode=0x1d f5=0x10 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] + data | at |
atom.ubase.and.b64 | M op opcode=0x1d f5=0x11 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] & data | at |
atom.ubase.or.b64 | M op opcode=0x1d f5=0x12 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] | data | at |
atom.ubase.xor.b64 | M op opcode=0x1d f5=0x13 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] ^ data | at |
atom.ubase.min.i64 | M op opcode=0x1d f5=0x14 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data) | at |
atom.ubase.max.i64 | M op opcode=0x1d f5=0x15 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data) | at |
atom.ubase.min.u64 | M op opcode=0x1d f5=0x16 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data) | at |
atom.ubase.max.u64 | M op opcode=0x1d f5=0x17 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data) | at |
atom.ubase.xchg.b64 | M op opcode=0x1d f5=0x18 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = data（交换） | at |
atom.ubase.cmpxchg.b64 | M op opcode=0x1d f5=0x19 sp=0x1 f2=0x0 | async:vst | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换） | at |
atom.buf.add.u32 | M op opcode=0x1d f5=0x0 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] + data | at bc |
atom.buf.sub.u32 | M op opcode=0x1d f5=0x1 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] - data | at bc |
atom.buf.and.b32 | M op opcode=0x1d f5=0x2 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] & data | at bc |
atom.buf.or.b32 | M op opcode=0x1d f5=0x3 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] | data | at bc |
atom.buf.xor.b32 | M op opcode=0x1d f5=0x4 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] ^ data | at bc |
atom.buf.min.i32 | M op opcode=0x1d f5=0x5 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_s(M[A], data) | at bc |
atom.buf.max.i32 | M op opcode=0x1d f5=0x6 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_s(M[A], data) | at bc |
atom.buf.min.u32 | M op opcode=0x1d f5=0x7 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_u(M[A], data) | at bc |
atom.buf.max.u32 | M op opcode=0x1d f5=0x8 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_u(M[A], data) | at bc |
atom.buf.xchg.b32 | M op opcode=0x1d f5=0x9 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = data（交换） | at bc |
atom.buf.cmpxchg.b32 | M op opcode=0x1d f5=0xa sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换） | at bc |
atom.buf.inc.u32 | M op opcode=0x1d f5=0xb sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = (old >= data) ? 0 : old + 1（回绕递增） | at bc |
atom.buf.dec.u32 | M op opcode=0x1d f5=0xc sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减） | at bc |
atom.buf.fadd.f32 | M op opcode=0x1d f5=0xd sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] + data（binary32 原子加，固定 RNE） | at bc opt |
atom.buf.fmin.f32 | M op opcode=0x1d f5=0xe sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = minimumNumber(M[A], data) | at bc opt |
atom.buf.fmax.f32 | M op opcode=0x1d f5=0xf sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = maximumNumber(M[A], data) | at bc opt |
atom.buf.add.u64 | M op opcode=0x1d f5=0x10 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] + data | at bc |
atom.buf.and.b64 | M op opcode=0x1d f5=0x11 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] & data | at bc |
atom.buf.or.b64 | M op opcode=0x1d f5=0x12 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] | data | at bc |
atom.buf.xor.b64 | M op opcode=0x1d f5=0x13 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] ^ data | at bc |
atom.buf.min.i64 | M op opcode=0x1d f5=0x14 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_s(M[A], data) | at bc |
atom.buf.max.i64 | M op opcode=0x1d f5=0x15 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_s(M[A], data) | at bc |
atom.buf.min.u64 | M op opcode=0x1d f5=0x16 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_u(M[A], data) | at bc |
atom.buf.max.u64 | M op opcode=0x1d f5=0x17 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_u(M[A], data) | at bc |
atom.buf.xchg.b64 | M op opcode=0x1d f5=0x18 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = data（交换） | at bc |
atom.buf.cmpxchg.b64 | M op opcode=0x1d f5=0x19 sp=0x2 f2=0x0 | async:vst | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换） | at bc |
atom.shared.add.u32 | M op opcode=0x1d f5=0x0 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] + data | at |
atom.shared.sub.u32 | M op opcode=0x1d f5=0x1 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] - data | at |
atom.shared.and.b32 | M op opcode=0x1d f5=0x2 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] & data | at |
atom.shared.or.b32 | M op opcode=0x1d f5=0x3 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] | data | at |
atom.shared.xor.b32 | M op opcode=0x1d f5=0x4 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] ^ data | at |
atom.shared.min.i32 | M op opcode=0x1d f5=0x5 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_s(M[A], data) | at |
atom.shared.max.i32 | M op opcode=0x1d f5=0x6 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_s(M[A], data) | at |
atom.shared.min.u32 | M op opcode=0x1d f5=0x7 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_u(M[A], data) | at |
atom.shared.max.u32 | M op opcode=0x1d f5=0x8 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_u(M[A], data) | at |
atom.shared.xchg.b32 | M op opcode=0x1d f5=0x9 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = data（交换） | at |
atom.shared.cmpxchg.b32 | M op opcode=0x1d f5=0xa sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：if M[A] == data0: M[A] = data1（比较交换） | at |
atom.shared.inc.u32 | M op opcode=0x1d f5=0xb sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = (old >= data) ? 0 : old + 1（回绕递增） | at |
atom.shared.dec.u32 | M op opcode=0x1d f5=0xc sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减） | at |
atom.shared.fadd.f32 | M op opcode=0x1d f5=0xd sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] + data（binary32 原子加，固定 RNE） | at opt |
atom.shared.fmin.f32 | M op opcode=0x1d f5=0xe sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = minimumNumber(M[A], data) | at opt |
atom.shared.fmax.f32 | M op opcode=0x1d f5=0xf sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = maximumNumber(M[A], data) | at opt |
atom.shared.add.u64 | M op opcode=0x1d f5=0x10 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] + data | at |
atom.shared.and.b64 | M op opcode=0x1d f5=0x11 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] & data | at |
atom.shared.or.b64 | M op opcode=0x1d f5=0x12 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] | data | at |
atom.shared.xor.b64 | M op opcode=0x1d f5=0x13 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] ^ data | at |
atom.shared.min.i64 | M op opcode=0x1d f5=0x14 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_s(M[A], data) | at |
atom.shared.max.i64 | M op opcode=0x1d f5=0x15 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_s(M[A], data) | at |
atom.shared.min.u64 | M op opcode=0x1d f5=0x16 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_u(M[A], data) | at |
atom.shared.max.u64 | M op opcode=0x1d f5=0x17 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_u(M[A], data) | at |
atom.shared.xchg.b64 | M op opcode=0x1d f5=0x18 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = data（交换） | at |
atom.shared.cmpxchg.b64 | M op opcode=0x1d f5=0x19 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：if M[A] == data0: M[A] = data1（比较交换） | at |
原子（返回旧值 .rtn）（104 条）
ATOM
语法（Syntax）
atom.global.add.u32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.sub.u32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.and.b32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.or.b32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.xor.b32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.i32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.i32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.u32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.u32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.xchg.b32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.cmpxchg.b32.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.inc.u32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.dec.u32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.fadd.f32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.fmin.f32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.fmax.f32.rtn rd:v* rs1:vp^2 imm:s24 cpol:cpol
atom.global.add.u64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.and.b64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.or.b64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.xor.b64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.i64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.i64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.min.u64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.max.u64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.xchg.b64.rtn rd:v^2* rs1:vp^2 imm:s24 cpol:cpol
atom.global.cmpxchg.b64.rtn rd:v^4* rs1:vp^2 imm:s24 cpol:cpol
atom.ubase.add.u32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.sub.u32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.and.b32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.or.b32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xor.b32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.i32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.i32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.u32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.u32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xchg.b32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.cmpxchg.b32.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.inc.u32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.dec.u32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.fadd.f32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.fmin.f32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.fmax.f32.rtn rd:v* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.add.u64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.and.b64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.or.b64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xor.b64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.i64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.i64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.min.u64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.max.u64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.xchg.b64.rtn rd:v^2* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.ubase.cmpxchg.b64.rtn rd:v^4* rs2:up^2 rs1:v imm:s24 cpol:cpol
atom.buf.add.u32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.sub.u32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.and.b32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.or.b32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xor.b32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.i32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.i32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.u32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.u32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xchg.b32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.cmpxchg.b32.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.inc.u32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.dec.u32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.fadd.f32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.fmin.f32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.fmax.f32.rtn rd:v* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.add.u64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.and.b64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.or.b64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xor.b64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.i64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.i64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.min.u64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.max.u64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.xchg.b64.rtn rd:v^2* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.buf.cmpxchg.b64.rtn rd:v^4* rs2:d4^4 rs1:v imm:s24 cpol:cpol
atom.shared.add.u32.rtn rd:v* rs1:v imm:s24
atom.shared.sub.u32.rtn rd:v* rs1:v imm:s24
atom.shared.and.b32.rtn rd:v* rs1:v imm:s24
atom.shared.or.b32.rtn rd:v* rs1:v imm:s24
atom.shared.xor.b32.rtn rd:v* rs1:v imm:s24
atom.shared.min.i32.rtn rd:v* rs1:v imm:s24
atom.shared.max.i32.rtn rd:v* rs1:v imm:s24
atom.shared.min.u32.rtn rd:v* rs1:v imm:s24
atom.shared.max.u32.rtn rd:v* rs1:v imm:s24
atom.shared.xchg.b32.rtn rd:v* rs1:v imm:s24
atom.shared.cmpxchg.b32.rtn rd:v^2* rs1:v imm:s24
atom.shared.inc.u32.rtn rd:v* rs1:v imm:s24
atom.shared.dec.u32.rtn rd:v* rs1:v imm:s24
atom.shared.fadd.f32.rtn rd:v* rs1:v imm:s24
atom.shared.fmin.f32.rtn rd:v* rs1:v imm:s24
atom.shared.fmax.f32.rtn rd:v* rs1:v imm:s24
atom.shared.add.u64.rtn rd:v^2* rs1:v imm:s24
atom.shared.and.b64.rtn rd:v^2* rs1:v imm:s24
atom.shared.or.b64.rtn rd:v^2* rs1:v imm:s24
atom.shared.xor.b64.rtn rd:v^2* rs1:v imm:s24
atom.shared.min.i64.rtn rd:v^2* rs1:v imm:s24
atom.shared.max.i64.rtn rd:v^2* rs1:v imm:s24
atom.shared.min.u64.rtn rd:v^2* rs1:v imm:s24
atom.shared.max.u64.rtn rd:v^2* rs1:v imm:s24
atom.shared.xchg.b64.rtn rd:v^2* rs1:v imm:s24
atom.shared.cmpxchg.b64.rtn rd:v^4* rs1:v imm:s24
| 变体 | 编码 | 延迟 | 语义 | 标志 |
|---|---|---|---|---|
atom.global.add.u32.rtn | M op opcode=0x21 f5=0x0 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] + data；旧值写回 rd | at |
atom.global.sub.u32.rtn | M op opcode=0x21 f5=0x1 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] - data；旧值写回 rd | at |
atom.global.and.b32.rtn | M op opcode=0x21 f5=0x2 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] & data；旧值写回 rd | at |
atom.global.or.b32.rtn | M op opcode=0x21 f5=0x3 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] | data；旧值写回 rd | at |
atom.global.xor.b32.rtn | M op opcode=0x21 f5=0x4 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] ^ data；旧值写回 rd | at |
atom.global.min.i32.rtn | M op opcode=0x21 f5=0x5 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data)；旧值写回 rd | at |
atom.global.max.i32.rtn | M op opcode=0x21 f5=0x6 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data)；旧值写回 rd | at |
atom.global.min.u32.rtn | M op opcode=0x21 f5=0x7 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data)；旧值写回 rd | at |
atom.global.max.u32.rtn | M op opcode=0x21 f5=0x8 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data)；旧值写回 rd | at |
atom.global.xchg.b32.rtn | M op opcode=0x21 f5=0x9 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = data（交换）；旧值写回 rd | at |
atom.global.cmpxchg.b32.rtn | M op opcode=0x21 f5=0xa sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at |
atom.global.inc.u32.rtn | M op opcode=0x21 f5=0xb sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = (old >= data) ? 0 : old + 1（回绕递增）；旧值写回 rd | at |
atom.global.dec.u32.rtn | M op opcode=0x21 f5=0xc sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减）；旧值写回 rd | at |
atom.global.fadd.f32.rtn | M op opcode=0x21 f5=0xd sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] + data（binary32 原子加，固定 RNE）；旧值写回 rd | at opt |
atom.global.fmin.f32.rtn | M op opcode=0x21 f5=0xe sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = minimumNumber(M[A], data)；旧值写回 rd | at opt |
atom.global.fmax.f32.rtn | M op opcode=0x21 f5=0xf sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = maximumNumber(M[A], data)；旧值写回 rd | at opt |
atom.global.add.u64.rtn | M op opcode=0x21 f5=0x10 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] + data；旧值写回 rd | at |
atom.global.and.b64.rtn | M op opcode=0x21 f5=0x11 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] & data；旧值写回 rd | at |
atom.global.or.b64.rtn | M op opcode=0x21 f5=0x12 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] | data；旧值写回 rd | at |
atom.global.xor.b64.rtn | M op opcode=0x21 f5=0x13 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = M[A] ^ data；旧值写回 rd | at |
atom.global.min.i64.rtn | M op opcode=0x21 f5=0x14 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data)；旧值写回 rd | at |
atom.global.max.i64.rtn | M op opcode=0x21 f5=0x15 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data)；旧值写回 rd | at |
atom.global.min.u64.rtn | M op opcode=0x21 f5=0x16 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data)；旧值写回 rd | at |
atom.global.max.u64.rtn | M op opcode=0x21 f5=0x17 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data)；旧值写回 rd | at |
atom.global.xchg.b64.rtn | M op opcode=0x21 f5=0x18 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：M[A] = data（交换）；旧值写回 rd | at |
atom.global.cmpxchg.b64.rtn | M op opcode=0x21 f5=0x19 sp=0x0 f2=0x0 | async:vm | A = s1_64 + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at |
atom.ubase.add.u32.rtn | M op opcode=0x21 f5=0x0 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] + data；旧值写回 rd | at |
atom.ubase.sub.u32.rtn | M op opcode=0x21 f5=0x1 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] - data；旧值写回 rd | at |
atom.ubase.and.b32.rtn | M op opcode=0x21 f5=0x2 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] & data；旧值写回 rd | at |
atom.ubase.or.b32.rtn | M op opcode=0x21 f5=0x3 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] | data；旧值写回 rd | at |
atom.ubase.xor.b32.rtn | M op opcode=0x21 f5=0x4 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] ^ data；旧值写回 rd | at |
atom.ubase.min.i32.rtn | M op opcode=0x21 f5=0x5 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data)；旧值写回 rd | at |
atom.ubase.max.i32.rtn | M op opcode=0x21 f5=0x6 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data)；旧值写回 rd | at |
atom.ubase.min.u32.rtn | M op opcode=0x21 f5=0x7 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data)；旧值写回 rd | at |
atom.ubase.max.u32.rtn | M op opcode=0x21 f5=0x8 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data)；旧值写回 rd | at |
atom.ubase.xchg.b32.rtn | M op opcode=0x21 f5=0x9 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = data（交换）；旧值写回 rd | at |
atom.ubase.cmpxchg.b32.rtn | M op opcode=0x21 f5=0xa sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at |
atom.ubase.inc.u32.rtn | M op opcode=0x21 f5=0xb sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = (old >= data) ? 0 : old + 1（回绕递增）；旧值写回 rd | at |
atom.ubase.dec.u32.rtn | M op opcode=0x21 f5=0xc sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减）；旧值写回 rd | at |
atom.ubase.fadd.f32.rtn | M op opcode=0x21 f5=0xd sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] + data（binary32 原子加，固定 RNE）；旧值写回 rd | at opt |
atom.ubase.fmin.f32.rtn | M op opcode=0x21 f5=0xe sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = minimumNumber(M[A], data)；旧值写回 rd | at opt |
atom.ubase.fmax.f32.rtn | M op opcode=0x21 f5=0xf sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = maximumNumber(M[A], data)；旧值写回 rd | at opt |
atom.ubase.add.u64.rtn | M op opcode=0x21 f5=0x10 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] + data；旧值写回 rd | at |
atom.ubase.and.b64.rtn | M op opcode=0x21 f5=0x11 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] & data；旧值写回 rd | at |
atom.ubase.or.b64.rtn | M op opcode=0x21 f5=0x12 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] | data；旧值写回 rd | at |
atom.ubase.xor.b64.rtn | M op opcode=0x21 f5=0x13 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = M[A] ^ data；旧值写回 rd | at |
atom.ubase.min.i64.rtn | M op opcode=0x21 f5=0x14 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_s(M[A], data)；旧值写回 rd | at |
atom.ubase.max.i64.rtn | M op opcode=0x21 f5=0x15 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_s(M[A], data)；旧值写回 rd | at |
atom.ubase.min.u64.rtn | M op opcode=0x21 f5=0x16 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = min_u(M[A], data)；旧值写回 rd | at |
atom.ubase.max.u64.rtn | M op opcode=0x21 f5=0x17 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = max_u(M[A], data)；旧值写回 rd | at |
atom.ubase.xchg.b64.rtn | M op opcode=0x21 f5=0x18 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：M[A] = data（交换）；旧值写回 rd | at |
atom.ubase.cmpxchg.b64.rtn | M op opcode=0x21 f5=0x19 sp=0x1 f2=0x0 | async:vm | A = U64[u2] + zext(s1) + sext(imm)；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at |
atom.buf.add.u32.rtn | M op opcode=0x21 f5=0x0 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] + data；旧值写回 rd | at bc |
atom.buf.sub.u32.rtn | M op opcode=0x21 f5=0x1 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] - data；旧值写回 rd | at bc |
atom.buf.and.b32.rtn | M op opcode=0x21 f5=0x2 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] & data；旧值写回 rd | at bc |
atom.buf.or.b32.rtn | M op opcode=0x21 f5=0x3 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] | data；旧值写回 rd | at bc |
atom.buf.xor.b32.rtn | M op opcode=0x21 f5=0x4 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] ^ data；旧值写回 rd | at bc |
atom.buf.min.i32.rtn | M op opcode=0x21 f5=0x5 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_s(M[A], data)；旧值写回 rd | at bc |
atom.buf.max.i32.rtn | M op opcode=0x21 f5=0x6 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_s(M[A], data)；旧值写回 rd | at bc |
atom.buf.min.u32.rtn | M op opcode=0x21 f5=0x7 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_u(M[A], data)；旧值写回 rd | at bc |
atom.buf.max.u32.rtn | M op opcode=0x21 f5=0x8 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_u(M[A], data)；旧值写回 rd | at bc |
atom.buf.xchg.b32.rtn | M op opcode=0x21 f5=0x9 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = data（交换）；旧值写回 rd | at bc |
atom.buf.cmpxchg.b32.rtn | M op opcode=0x21 f5=0xa sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at bc |
atom.buf.inc.u32.rtn | M op opcode=0x21 f5=0xb sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = (old >= data) ? 0 : old + 1（回绕递增）；旧值写回 rd | at bc |
atom.buf.dec.u32.rtn | M op opcode=0x21 f5=0xc sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减）；旧值写回 rd | at bc |
atom.buf.fadd.f32.rtn | M op opcode=0x21 f5=0xd sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] + data（binary32 原子加，固定 RNE）；旧值写回 rd | at bc opt |
atom.buf.fmin.f32.rtn | M op opcode=0x21 f5=0xe sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = minimumNumber(M[A], data)；旧值写回 rd | at bc opt |
atom.buf.fmax.f32.rtn | M op opcode=0x21 f5=0xf sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = maximumNumber(M[A], data)；旧值写回 rd | at bc opt |
atom.buf.add.u64.rtn | M op opcode=0x21 f5=0x10 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] + data；旧值写回 rd | at bc |
atom.buf.and.b64.rtn | M op opcode=0x21 f5=0x11 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] & data；旧值写回 rd | at bc |
atom.buf.or.b64.rtn | M op opcode=0x21 f5=0x12 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] | data；旧值写回 rd | at bc |
atom.buf.xor.b64.rtn | M op opcode=0x21 f5=0x13 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = M[A] ^ data；旧值写回 rd | at bc |
atom.buf.min.i64.rtn | M op opcode=0x21 f5=0x14 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_s(M[A], data)；旧值写回 rd | at bc |
atom.buf.max.i64.rtn | M op opcode=0x21 f5=0x15 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_s(M[A], data)；旧值写回 rd | at bc |
atom.buf.min.u64.rtn | M op opcode=0x21 f5=0x16 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = min_u(M[A], data)；旧值写回 rd | at bc |
atom.buf.max.u64.rtn | M op opcode=0x21 f5=0x17 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = max_u(M[A], data)；旧值写回 rd | at bc |
atom.buf.xchg.b64.rtn | M op opcode=0x21 f5=0x18 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：M[A] = data（交换）；旧值写回 rd | at bc |
atom.buf.cmpxchg.b64.rtn | M op opcode=0x21 f5=0x19 sp=0x2 f2=0x0 | async:vm | A = DESC(u2..u2+3).base + s1 + sext(imm)（按 DESC.size 越界检查）；（L2 原子）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at bc |
atom.shared.add.u32.rtn | M op opcode=0x21 f5=0x0 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] + data；旧值写回 rd | at |
atom.shared.sub.u32.rtn | M op opcode=0x21 f5=0x1 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] - data；旧值写回 rd | at |
atom.shared.and.b32.rtn | M op opcode=0x21 f5=0x2 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] & data；旧值写回 rd | at |
atom.shared.or.b32.rtn | M op opcode=0x21 f5=0x3 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] | data；旧值写回 rd | at |
atom.shared.xor.b32.rtn | M op opcode=0x21 f5=0x4 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] ^ data；旧值写回 rd | at |
atom.shared.min.i32.rtn | M op opcode=0x21 f5=0x5 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_s(M[A], data)；旧值写回 rd | at |
atom.shared.max.i32.rtn | M op opcode=0x21 f5=0x6 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_s(M[A], data)；旧值写回 rd | at |
atom.shared.min.u32.rtn | M op opcode=0x21 f5=0x7 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_u(M[A], data)；旧值写回 rd | at |
atom.shared.max.u32.rtn | M op opcode=0x21 f5=0x8 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_u(M[A], data)；旧值写回 rd | at |
atom.shared.xchg.b32.rtn | M op opcode=0x21 f5=0x9 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = data（交换）；旧值写回 rd | at |
atom.shared.cmpxchg.b32.rtn | M op opcode=0x21 f5=0xa sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at |
atom.shared.inc.u32.rtn | M op opcode=0x21 f5=0xb sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = (old >= data) ? 0 : old + 1（回绕递增）；旧值写回 rd | at |
atom.shared.dec.u32.rtn | M op opcode=0x21 f5=0xc sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = (old == 0 || old > data) ? data : old - 1（回绕递减）；旧值写回 rd | at |
atom.shared.fadd.f32.rtn | M op opcode=0x21 f5=0xd sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] + data（binary32 原子加，固定 RNE）；旧值写回 rd | at opt |
atom.shared.fmin.f32.rtn | M op opcode=0x21 f5=0xe sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = minimumNumber(M[A], data)；旧值写回 rd | at opt |
atom.shared.fmax.f32.rtn | M op opcode=0x21 f5=0xf sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = maximumNumber(M[A], data)；旧值写回 rd | at opt |
atom.shared.add.u64.rtn | M op opcode=0x21 f5=0x10 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] + data；旧值写回 rd | at |
atom.shared.and.b64.rtn | M op opcode=0x21 f5=0x11 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] & data；旧值写回 rd | at |
atom.shared.or.b64.rtn | M op opcode=0x21 f5=0x12 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] | data；旧值写回 rd | at |
atom.shared.xor.b64.rtn | M op opcode=0x21 f5=0x13 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = M[A] ^ data；旧值写回 rd | at |
atom.shared.min.i64.rtn | M op opcode=0x21 f5=0x14 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_s(M[A], data)；旧值写回 rd | at |
atom.shared.max.i64.rtn | M op opcode=0x21 f5=0x15 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_s(M[A], data)；旧值写回 rd | at |
atom.shared.min.u64.rtn | M op opcode=0x21 f5=0x16 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = min_u(M[A], data)；旧值写回 rd | at |
atom.shared.max.u64.rtn | M op opcode=0x21 f5=0x17 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = max_u(M[A], data)；旧值写回 rd | at |
atom.shared.xchg.b64.rtn | M op opcode=0x21 f5=0x18 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：M[A] = data（交换）；旧值写回 rd | at |
atom.shared.cmpxchg.b64.rtn | M op opcode=0x21 f5=0x19 sp=0x3 f2=0x0 | async:ls | A = s1 + sext(imm)（workgroup shared 窗口）；（shared 内）：if M[A] == data0: M[A] = data1（比较交换）；旧值写回 rd | at |
uniform 标量加载（20 条）
ULD
语法（Syntax）
uld.b32 rd:u rs1:up^2 imm:s24 cpol:cpol
uld.b64 rd:u^2 rs1:up^2 imm:s24 cpol:cpol
uld.b128 rd:u^4 rs1:up^2 imm:s24 cpol:cpol
uld.b256 rd:u^8 rs1:up^2 imm:s24 cpol:cpol
uld.b512 rd:u^16 rs1:up^2 imm:s24 cpol:cpol
uld.x.b32 rd:u rs1:up^2 rs2:u imm:s24 cpol:cpol
uld.x.b64 rd:u^2 rs1:up^2 rs2:u imm:s24 cpol:cpol
uld.x.b128 rd:u^4 rs1:up^2 rs2:u imm:s24 cpol:cpol
uld.x.b256 rd:u^8 rs1:up^2 rs2:u imm:s24 cpol:cpol
uld.x.b512 rd:u^16 rs1:up^2 rs2:u imm:s24 cpol:cpol
uld.buf.b32 rd:u rs1:d4^4 imm:s24 cpol:cpol
uld.buf.b64 rd:u^2 rs1:d4^4 imm:s24 cpol:cpol
uld.buf.b128 rd:u^4 rs1:d4^4 imm:s24 cpol:cpol
uld.buf.b256 rd:u^8 rs1:d4^4 imm:s24 cpol:cpol
uld.buf.b512 rd:u^16 rs1:d4^4 imm:s24 cpol:cpol
uld.bufx.b32 rd:u rs1:d4^4 rs2:u imm:s24 cpol:cpol
uld.bufx.b64 rd:u^2 rs1:d4^4 rs2:u imm:s24 cpol:cpol
uld.bufx.b128 rd:u^4 rs1:d4^4 rs2:u imm:s24 cpol:cpol
uld.bufx.b256 rd:u^8 rs1:d4^4 rs2:u imm:s24 cpol:cpol
uld.bufx.b512 rd:u^16 rs1:d4^4 rs2:u imm:s24 cpol:cpol
| 变体 | 编码 | 延迟 | 语义 | 标志 |
|---|---|---|---|---|
uld.b32 | UM op opcode=0x25 f7=0x0 | async:sm | A = U64[u1] + sext(imm)；加载 1 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.b64 | UM op opcode=0x25 f7=0x1 | async:sm | A = U64[u1] + sext(imm)；加载 2 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.b128 | UM op opcode=0x25 f7=0x2 | async:sm | A = U64[u1] + sext(imm)；加载 4 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.b256 | UM op opcode=0x25 f7=0x3 | async:sm | A = U64[u1] + sext(imm)；加载 8 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.b512 | UM op opcode=0x25 f7=0x4 | async:sm | A = U64[u1] + sext(imm)；加载 16 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.x.b32 | UM op opcode=0x25 f7=0x8 | async:sm | A = U64[u1] + u2 + sext(imm)；加载 1 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.x.b64 | UM op opcode=0x25 f7=0x9 | async:sm | A = U64[u1] + u2 + sext(imm)；加载 2 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.x.b128 | UM op opcode=0x25 f7=0xa | async:sm | A = U64[u1] + u2 + sext(imm)；加载 4 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.x.b256 | UM op opcode=0x25 f7=0xb | async:sm | A = U64[u1] + u2 + sext(imm)；加载 8 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.x.b512 | UM op opcode=0x25 f7=0xc | async:sm | A = U64[u1] + u2 + sext(imm)；加载 16 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei |
uld.buf.b32 | UM op opcode=0x25 f7=0x10 | async:sm | A = DESC(u1..u1+3).base + sext(imm)（越界 / null 读 0）；加载 1 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.buf.b64 | UM op opcode=0x25 f7=0x11 | async:sm | A = DESC(u1..u1+3).base + sext(imm)（越界 / null 读 0）；加载 2 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.buf.b128 | UM op opcode=0x25 f7=0x12 | async:sm | A = DESC(u1..u1+3).base + sext(imm)（越界 / null 读 0）；加载 4 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.buf.b256 | UM op opcode=0x25 f7=0x13 | async:sm | A = DESC(u1..u1+3).base + sext(imm)（越界 / null 读 0）；加载 8 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.buf.b512 | UM op opcode=0x25 f7=0x14 | async:sm | A = DESC(u1..u1+3).base + sext(imm)（越界 / null 读 0）；加载 16 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.bufx.b32 | UM op opcode=0x25 f7=0x18 | async:sm | A = DESC(u1..u1+3).base + u2 + sext(imm)（越界读 0）；加载 1 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.bufx.b64 | UM op opcode=0x25 f7=0x19 | async:sm | A = DESC(u1..u1+3).base + u2 + sext(imm)（越界读 0）；加载 2 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.bufx.b128 | UM op opcode=0x25 f7=0x1a | async:sm | A = DESC(u1..u1+3).base + u2 + sext(imm)（越界读 0）；加载 4 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.bufx.b256 | UM op opcode=0x25 f7=0x1b | async:sm | A = DESC(u1..u1+3).base + u2 + sext(imm)（越界读 0）；加载 8 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |
uld.bufx.b512 | UM op opcode=0x25 f7=0x1c | async:sm | A = DESC(u1..u1+3).base + u2 + sext(imm)（越界读 0）；加载 16 个 dword 到 U[d..]（无视 EXEC，标量路径） | ld ei bc |