WAVE specification v0.2

• Version: 0.2, working draft
• Authors: Ojima Abraham, Onyinye Okoli
• Date: March 23, 2026
• Status: working draft

0. Conformance language

The terms MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL follow RFC 2119.

1. Introduction

1.1 Purpose

WAVE means Wide Architecture Virtual Encoding. This specification defines a vendor-neutral instruction set for GPU computation. It covers execution, registers, memory, control flow, instructions, and capability queries.

The thin abstraction principle defines what a compliant implementation MUST do while leaving the mechanism open. Implementations MAY use any microarchitectural technique that meets these requirements.

1.2 Scope

This specification covers compute workloads. It excludes rasterization, tessellation, pixel export, and ray tracing. Future extensions MAY add those operations.

1.3 Design principles

1. Thin abstraction. Every requirement traces to a hardware-invariant primitive observed across all four major GPU vendors. No requirement is imposed for software convenience alone.
2. Queryable parameters. Programs query wave width, register count, and scratchpad size. These values can differ across implementations.
3. Structured divergence. The specification defines control flow semantics but not divergence mechanisms. Implementations can use execution masks, predication, hardware stacks, or per-thread program counters.
4. Mandatory minimums. Every queryable parameter has a minimum value. A compliant implementation MUST meet or exceed all minimums.

1.4 Relationship to other standards

SPIR-V MAY be used as a distribution format for programs targeting this ISA. OpenCL and Vulkan MAY serve as host APIs for dispatching workloads. This specification defines the hardware execution model; existing standards define host-device interaction.

1.5 Encoding and execution

1. Register fields are 8 bits and address up to 256 registers. See Section 8.2.
2. Implementations MUST support at least 32 levels of nested control flow. See Sections 5.4 and 7.1.
3. Predicate negation applies to break and continue. See Section 5.1.
4. Each wave MUST have independent control flow state. See Section 5.5.
5. Appendix A lists opcodes. Section 9.4 describes conformance tests.

2. Execution model

2.1 Overview

A compliant processor has one or more cores, each able to execute multiple workgroups concurrently. Software cannot address cores. Hardware assigns workgroups to cores, and the programmer MUST NOT assume any particular mapping.

2.2 Thread hierarchy

The execution model defines four levels, three mandatory and one optional:

Level 0: thread. The smallest unit of execution. A thread has a private register file, a scalar program counter, and hardware-set identity values for its position in the hierarchy. A thread executes a sequential stream of instructions.

Level 1: wave. A group of exactly W threads that execute a single instruction simultaneously, where W is a hardware constant queryable at compile time. See Section 7. All threads in a wave share a program counter for instruction fetch. When threads in a wave disagree on a branch condition, the implementation MUST ensure that both paths execute with inactive threads producing no architectural side effects. The implementation chooses the mechanism.

A wave is the fundamental scheduling unit. The hardware scheduler operates on waves, not individual threads. Each wave MUST maintain independent control flow state. See Section 5.5.

Level 2: workgroup. A group of one or more waves, containing up to MAX_WORKGROUP_SIZE threads. All waves in a workgroup execute on the same core, share access to local memory, and may synchronize via barriers.

Waves per workgroup: ceil(workgroup_thread_count / W).

Workgroup dimensions are specified at dispatch time as x, y, and z, where x * y * z <= MAX_WORKGROUP_SIZE.

Level 3: grid. The complete dispatch of workgroups. A grid is specified as a 3-dimensional count of workgroups. Workgroups within a grid MAY execute in any order, on any core, at any time. No synchronization is available between workgroups within a single grid dispatch.

Level 2.5: optional cluster. A group of workgroups guaranteed to execute concurrently on adjacent cores with access to each other's local memory. The cluster size is queryable via CLUSTER_SIZE. If the implementation does not support clusters, CLUSTER_SIZE is 1 and cluster-scope operations behave identically to workgroup-scope operations.

2.3 Thread identifiers

Every thread exposes these hardware-populated, read-only special registers:

Identifier	Type	Description
thread_id.{x,y,z}	uint32	Thread position within workgroup; 3D
wave_id	uint32	Wave index within workgroup
lane_id	uint32	Thread position within wave; 0 to W-1
workgroup_id.{x,y,z}	uint32	Workgroup position within grid; 3D
workgroup_size.{x,y,z}	uint32	Workgroup dimensions
grid_size.{x,y,z}	uint32	Grid dimensions; in workgroups
num_waves	uint32	Number of waves in this workgroup

2.4 Core resources

Each core provides these resources:

The register file has a fixed capacity of F bytes of on-chip storage, partitioned among resident threads. Each thread receives R registers, declared at compile time. The number of resident waves is bounded by:

max_resident_waves = floor(F / (R * W * 4))

where each register occupies 4 bytes, or 32 bits. Implementations MUST support at least MIN_MAX_REGISTERS registers per thread.

Local memory is a fixed-size on-chip scratchpad of S bytes shared by all waves in a workgroup. Load and store instructions address it explicitly; there is no automatic caching or data placement. Its contents are undefined at workgroup start and do not survive workgroup boundaries. Implementations MUST provide at least MIN_LOCAL_MEMORY_SIZE bytes.

The hardware scheduler selects a ready wave for execution each cycle. When a wave stalls on a memory access, barrier, or other long-latency operation, the scheduler MUST be able to select another resident wave without software-visible overhead. The scheduling policy is implementation-defined.

2.5 Execution guarantees

1. All threads in a wave execute the same instruction in the same cycle, or appear to from the programmer's perspective.
2. Within a wave, instruction execution is in program order.
3. Between waves in the same workgroup, no execution order is guaranteed unless explicitly synchronized via barriers or memory ordering operations.
4. Between workgroups, no execution order is guaranteed.
5. A workgroup that uses local memory or barriers MUST have all its waves resident on a single core simultaneously.
6. The implementation MUST guarantee forward progress for at least one wave per core at all times. Scheduling cannot cause deadlock.
7. A wave that has been dispatched MUST eventually complete, assuming the program terminates. Scheduling cannot starve a wave.
8. Execution MUST be deterministic: given identical inputs, program binary, and dispatch configuration, the same program MUST produce identical results across runs on the same implementation. Non-determinism across different implementations is permitted for implementation-defined behaviors.

2.6 Dispatch

A dispatch launches a grid of workgroups and specifies:

• The kernel entry point
• Grid dimensions
• Workgroup dimensions
• Register count R per thread
• Local memory size required, in bytes
• Kernel arguments: buffer addresses and constants

The implementation MUST reject a dispatch if the requested resources exceed core capacity.

3. Register model

3.1 General-purpose registers

Each thread has access to R general-purpose registers, where R is declared at compile time and MUST NOT exceed MAX_REGISTERS. Registers are 32 bits wide. They are untyped at the hardware level; the instruction determines whether the register holds an integer, float, or bitfield. Registers are named r0 through r{R-1}.

3.2 Sub-register access

Each 32-bit register MAY be accessed as two 16-bit halves: r{N}.lo, bits [15:0], and r{N}.hi, bits [31:16]. This enables efficient F16 and BF16 operations without consuming additional registers.

3.3 Register pairs

Two consecutive registers MAY hold a 64-bit value: r{N}:r{N+1} with r{N} as the low 32 bits. Register pairs support 64-bit integer operations and F64 operations where available.

3.4 Special registers

Hardware-populated read-only registers for thread identity:

• sr_thread_id_x, sr_thread_id_y, sr_thread_id_z
• sr_wave_id, sr_lane_id
• sr_workgroup_id_x, sr_workgroup_id_y, sr_workgroup_id_z
• sr_workgroup_size_x, sr_workgroup_size_y, sr_workgroup_size_z
• sr_grid_size_x, sr_grid_size_y, sr_grid_size_z
• sr_wave_width, sr_num_waves

3.5 Predicate registers

The implementation MUST provide at least 4 predicate registers, p0 through p3, each 1 bit wide per thread. Predicates are set by comparison instructions and consumed by conditional branch instructions.

3.6 Allocation and occupancy

The register count R is declared for each kernel at compile time. Every thread in each resident wave receives R registers. The occupancy equation determines how many waves can reside simultaneously. Compilers SHOULD minimize R to maximize occupancy.

4. Memory model

4.1 Memory spaces

Three mandatory memory spaces:

• Register memory: per-thread, on-chip, single-cycle, not addressable
• Local memory: per-workgroup, on-chip, explicitly addressed, size S
• Device memory: global, off-chip or unified, cached by hardware-managed caches, persistent across dispatches

4.2 Local memory details

Local memory is organized as a flat byte-addressable array of S bytes. The base address is 0. Addresses outside [0, S) produce undefined behavior.

Local memory is banked. When multiple threads in a wave access the same bank in the same cycle, a bank conflict MAY occur. Broadcasts to the same address within a bank MUST NOT cause a conflict.

Access widths are 8, 16, 32, and 64 bits.

4.3 Device memory details

Device memory is byte-addressable with a 64-bit virtual address space. The implementation MUST support aligned loads and stores of 8, 16, 32, 64, and 128 bits. Coalescing of contiguous accesses is implementation-defined. Cache hierarchy is transparent to the ISA.

4.4 Memory ordering

Default ordering is relaxed. Fence operations provide ordering at four scopes:

• scope_wave
• scope_workgroup
• scope_device
• scope_system

Fence semantics:

• fence_acquire(scope) ensures subsequent loads see values at least as recent as those visible at the scope
• fence_release(scope) ensures prior stores are visible at the scope
• fence_acq_rel(scope) combines both

Store-to-load ordering within a thread is always guaranteed.

4.5 Atomic operations

Atomic operations perform indivisible read-modify-write sequences on local memory and device memory. Required operations:

• atomic_add, i32, u32, f32
• atomic_sub, i32, u32
• atomic_min, i32, u32
• atomic_max, i32, u32
• atomic_and, u32
• atomic_or, u32
• atomic_xor, u32
• atomic_exchange, u32
• atomic_compare_swap, u32

Each takes a scope parameter. 64-bit atomics are OPTIONAL.

5. Control flow

5.1 Structured control flow

The ISA defines structured control flow primitives. All control flow MUST be expressible through these primitives. The implementation MUST NOT require the programmer to manage divergence masks, execution masks, or reconvergence points.

Conditional:

if (predicate)
    <then-body>
else
    <else-body>
endif

When threads in a wave evaluate the predicate differently, the implementation MUST execute both paths. Threads for which the predicate is false MUST NOT produce side effects during the then-body, and vice versa for the else-body. After endif, all threads that were active before the if are active again.

Loop:

loop
    <body>
    break (predicate)    // exit loop for Threads where predicate is true
    continue (predicate) // skip to next iteration for Threads where predicate is true
endloop

A loop executes until all active threads have exited via break. The implementation MUST guarantee forward progress: if at least one thread remains in the loop, execution continues.

Break and continue with a negated predicate, such as break !p0, apply to threads where the predicate is false. The instruction break !p0 exits the loop for threads where p0 is false. Negation is applied before selecting affected threads.

Function call:

call <function>
return

Function calls push the return address onto an implementation-managed stack that MUST support at least MAX_CALL_DEPTH nested calls. See Section 7.1. Recursion is OPTIONAL. See Section 7.

5.2 Uniform branches

If all threads in a wave evaluate a branch identically, the implementation SHOULD avoid executing the not-taken path.

5.3 Divergence and reconvergence

Implementations can use any mechanism that preserves the structured control flow semantics of Section 5.1:

• Compiler-managed execution masks, AMD approach
• Hardware per-thread program counters, NVIDIA approach
• Compiler-generated predicated instructions, Intel approach
• Hardware divergence stack, Apple approach
• Any other mechanism that preserves the specified semantics

The ISA does not expose or constrain the divergence mechanism.

5.4 Minimum divergence depth

Implementations MUST support nested if/else/endif and loop/break/endloop constructs to at least MIN_DIVERGENCE_DEPTH levels. See Section 7.1. The limit applies to any mix of nested conditionals and loops.

If a program exceeds the implementation's maximum divergence depth, the behavior is undefined.

5.5 Per-wave control flow state

Each wave MUST maintain independent control flow state. This state includes the divergence stack, active mask history, loop iteration state, and reconvergence points. Two waves executing the same program binary at different points in a loop or branch MUST NOT interfere with each other's control flow state.

Independent control flow state lets waves reach barriers at different loop iterations without corrupting each other's masks.

6. Instruction set

6.1 Instruction format

All instructions operate on registers. There are no memory-to-register or memory-to-memory instructions, except explicit load/store.

Assembly notation:

opcode destination, source1, source2

Predicated instructions are written as:

@predicate opcode destination, source1, source2     // execute if predicate is true
@!predicate opcode destination, source1, source2    // execute if predicate is false

A predicated instruction executes only for threads whose predicate condition is met. Other threads retain their destination register values and produce no side effects, including memory stores and atomics.

The binary encoding of instructions is defined in Section 8.

6.2 Integer arithmetic

Integer operations act per thread on i32/u32 values:

• iadd, isub, imul, imul_hi, imad
• idiv, imod, ineg, iabs
• imin, imax, iclamp

Integer arithmetic uses wrapping semantics for overflow. Division by zero produces undefined behavior.

6.3 Bitwise operations

• and, or, xor, not
• shl, shr, logical, sar, arithmetic
• bitcount, bitfind, bitrev
• bfe, bit field extract, bfi, bit field insert

Shift amounts are masked to 5 bits, shift by rs2 & 0x1F.

6.4 Floating-point arithmetic: F32: REQUIRED

IEEE 754 single precision. Operations:

• fadd, fsub, fmul, fma, fdiv
• fneg, fabs, fmin, fmax, fclamp
• fsqrt, frsqrt, frcp
• ffloor, fceil, fround, ftrunc, ffract, fsat

Transcendentals: fsin, fcos, fexp2, flog2, at least 2 ULP precision. Denormals MAY be flushed to zero.

6.5 Floating-point arithmetic: F16: REQUIRED

F16 operations on register halves. Packed 2xF16 operations for throughput:

• hadd2, hmul2, hma2

6.6 Floating-point arithmetic: F64: OPTIONAL

F64 operations on register pairs:

• dadd, dsub, dmul, dma, ddiv, dsqrt

6.7 Type conversion

Type conversion between i32, u32, f16, f32, f64.

6.8 Comparison and select

Comparison instructions set predicate registers. Select instructions provide conditional moves.

6.9 Memory operations

Local memory: local_load/local_store for u8, u16, u32, u64.

Device memory: device_load/device_store for u8, u16, u32, u64, u128. Device loads are asynchronous; use wait or fence before consuming.

Optional cache hints: .cached, .uncached, .streaming.

6.10 Atomic operations

Atomic instructions operate on local and device memory with scope suffixes .wave, .workgroup, .device, .system. They return the old value; non-returning variants SHOULD be optimized.

6.11 Wave operations: REQUIRED

Wave operations communicate between threads within a wave without going through memory.

• wave_shuffle, by lane, wave_shuffle_up, wave_shuffle_down, wave_shuffle_xor
• wave_broadcast
• wave_ballot, wave_any, wave_all
• wave_prefix_sum, exclusive
• wave_reduce_add, wave_reduce_min, wave_reduce_max

Shuffle results are implementation-defined when the source lane is inactive or outside the range 0 <= lane < W.

Wave operations use only active threads. Threads masked by divergence do not participate in reductions, ballots, or prefix sums. Their registers stay unchanged.

6.12 Synchronization

• barrier: workgroup-scope barrier. All waves in the workgroup MUST reach this point before any wave proceeds past it. Memory operations before the barrier are visible to all waves in the workgroup after the barrier.
• fence_acquire/fence_release/fence_acq_rel, scoped
• wait, for async loads

A barrier instruction MUST NOT appear inside divergent control flow. When a wave reaches a barrier, all threads in that wave that were active at the outermost non-divergent scope must reach the same barrier. Barriers are permitted inside uniform if blocks where all threads agree. Loops may contain barriers if all waves in the workgroup execute the same number of barrier instructions per loop iteration.

6.13 Control flow instructions

Instruction	Description
if pd	Begin conditional block; threads where pd is false become inactive
else	Switch active/inactive threads
endif	End conditional block; restore original active set
loop	Begin loop
break pd	Threads where pd is true exit the loop
break !pd	Threads where pd is false exit the loop
continue pd	Threads where pd is true skip to next iteration
continue !pd	Threads where pd is false skip to next iteration
endloop	End loop; branch back to loop if any threads still active
call <label>	Call function
return	Return from function
halt	Terminate this thread

6.14 Matrix multiply-accumulate: OPTIONAL

• mma_f16_f32, mma_bf16_f32, mma_f32_f32

Tile dimensions are queryable.

Miscellaneous

• mov, mov_imm, nop

7. Capability system

7.1 Required constants

Constant	Minimum	Description
WAVE_WIDTH	8	Threads per wave
MAX_REGISTERS	64	Maximum registers per thread
REGISTER_FILE_SIZE	16384	Total register file size; bytes
LOCAL_MEMORY_SIZE	16384	Local memory per workgroup; bytes
MAX_WORKGROUP_SIZE	256	Maximum threads per workgroup
MAX_WORKGROUPS_PER_CORE	1	Maximum concurrent workgroups per core
MAX_WAVES_PER_CORE	4	Maximum concurrent waves per core
DEVICE_MEMORY_SIZE	:	Total device memory; bytes
CLUSTER_SIZE	1	Workgroups per cluster
MAX_CALL_DEPTH	8	Maximum function call depth
MIN_DIVERGENCE_DEPTH	32	Minimum nested control flow depth

7.2 Optional capabilities

• CAP_F64: 64-bit floating point
• CAP_ATOMIC_64: 64-bit atomics
• CAP_ATOMIC_F32: F32 atomic add
• CAP_MMA: matrix multiply-accumulate
• CAP_RECURSION: recursive function calls
• CAP_CLUSTER: cluster support

7.3 Matrix MMA parameters

When CAP_MMA is present, the following are queryable:

• MMA_M, MMA_N, MMA_K: tile dimensions
• MMA_TYPES: supported input/output type combinations

7.4 Query mechanism

The host API provides query_constant and query_capability functions.

8. Binary encoding

8.1 Overview

Instructions use fixed-width 48-bit words of 6 bytes. Some instructions add a 32-bit word for immediate values or extended operands, giving a total of 80 bits / 10 bytes.

The v0.2 encoding uses 48 bits with 8-bit register fields for up to 256 registers. The v0.1 format uses 32-bit base instructions with 5-bit register fields for up to 32 registers, below MAX_REGISTERS = 64.

8.2 Base instruction format: 48-bit

Bits	Field	Description
[47:40]	opcode	8 bits: 256 primary opcodes
[39:32]	rd	8 bits: destination register; 0-255
[31:24]	rs1	8 bits: source register 1; 0-255
[23:16]	rs2	8 bits: source register 2; 0-255
[15:12]	modifier	4 bits: instruction-specific sub-opcode
[11:10]	scope	2 bits: memory scope; 00=wave, 01=workgroup, 10=device, 11=system
[9:8]	pred	2 bits: predicate register selector; p0-p3
[7]	pred_neg	1 bit: negate predicate; 0=normal, 1=negated
[6]	pred_en	1 bit: predication enable; 0=unpredicated, 1=predicated
[5:0]	flags	6 bits: instruction-specific

8.3 Extended instruction format: 80-bit

For instructions requiring a third source register or a 32-bit immediate:

• Word 0, 48 bits: Base instruction as above, with flags indicating extended format
• Word 1, 32 bits: [31:0] rs3, 8 bits, + imm24, or full imm32

8.4 Opcode map

The 8-bit opcode field provides 256 primary opcodes, organized as:

Range	Category
0x00-0x0F	Integer arithmetic
0x10-0x1F	Floating-point arithmetic; F32
0x20-0x27	Bitwise operations
0x28-0x2F	Comparison and select
0x30-0x37	Local memory operations
0x38-0x3F	Device memory operations
0x40-0x4F	Atomic operations
0x50-0x5F	Wave operations
0x60-0x6F	Control flow and synchronization
0x70-0x7F	Type conversion
0x80-0x8F	F16 arithmetic
0x90-0x9F	F64 arithmetic; optional
0xA0-0xAF	Matrix MMA; optional
0xB0-0xEF	Reserved for future extensions
0xF0-0xFF	Miscellaneous; mov, mov_imm, nop, halt

See Appendix A for the complete opcode-to-mnemonic mapping.

9. Conformance

9.1 Required behavior

A compliant implementation MUST:

1. Support all mandatory instructions, Sections 6.2 through 6.15.
2. Meet or exceed all minimum values in Section 7.1.
3. Implement the memory ordering semantics of Section 4.4.
4. Implement the structured control flow semantics of Section 5.1, including per-wave control flow state, Section 5.5.
5. Satisfy all execution guarantees of Section 2.5.
6. Correctly report all capabilities of Section 7.2.
7. Support nested control flow to at least MIN_DIVERGENCE_DEPTH levels, Section 5.4.

9.2 Implementation-defined behavior

Valid implementations may differ in these behaviors:

• Denormal floating-point handling, flush to zero or preserve
• Bank conflict penalty in local memory
• Device memory coalescing policy
• Cache hierarchy structure, sizes, and policies
• Scheduling policy for wave selection
• Transcendental function precision beyond the specified minimum
• Out-of-bounds shuffle source lane result
• Shuffle from inactive source lane result
• Unaligned memory access behavior
• Wave scheduling order within a workgroup

9.3 Undefined behavior

The following have undefined behavior:

• Accessing local memory outside [0, S)
• Accessing device memory outside allocated regions
• Using optional capabilities on hardware that does not support them
• Exceeding MAX_CALL_DEPTH
• Exceeding the implementation's maximum divergence depth
• Data races on device memory without proper fencing
• Infinite loops with no forward progress
• Barrier inside a divergent control flow path, where threads in a wave disagree on whether to execute the barrier
• Integer division by zero

9.4 Conformance testing

The WAVE toolchain repository includes a conformance test suite. It verifies:

1. Correct execution of all mandatory instructions, including overflow, NaN, and infinity
2. Memory ordering compliance across scopes
3. Barrier semantics with multi-wave workgroups, including barriers inside loops
4. Atomic operation correctness on both local and device memory
5. Structured control flow behavior under divergence, including nested divergence to depth 32
6. Wave operations under divergence, shuffle, ballot, reduce with inactive threads
7. Capability reporting accuracy
8. Real GPU program correctness, tiled GEMM, parallel reduction, histogram, prefix sum

An implementation passes conformance if all tests in the mandatory suite produce correct results. The test suite is versioned alongside the specification.

A. Full opcode table

Integer arithmetic: 0x00-0x0F

Opcode	Mnemonic	Format	Description
0x00	iadd	Base	rd = rs1 + rs2
0x01	isub	Base	rd = rs1 - rs2
0x02	imul	Base	rd = (rs1 * rs2) & 0xFFFFFFFF
0x03	imul_hi	Base	rd = (rs1 * rs2) >> 32
0x04	imad	Extended	rd = rs1 * rs2 + rs3
0x05	idiv	Base	rd = rs1 / rs2
0x06	imod	Base	rd = rs1 % rs2
0x07	ineg	Base	rd = -rs1
0x08	iabs	Base	rd = abs(rs1)
0x09	imin	Base	rd = min(rs1, rs2); signed
0x0A	imax	Base	rd = max(rs1, rs2); signed
0x0B	iclamp	Extended	rd = clamp(rs1, rs2, rs3)
0x0C	umin	Base	rd = min(rs1, rs2); unsigned
0x0D	umax	Base	rd = max(rs1, rs2); unsigned
0x0E-0x0F	Reserved	:	:

Floating-point arithmetic F32: 0x10-0x1F

Opcode	Mnemonic	Format	Description
0x10	fadd	Base	rd = rs1 + rs2
0x11	fsub	Base	rd = rs1 - rs2
0x12	fmul	Base	rd = rs1 * rs2
0x13	fma	Extended	rd = rs1 * rs2 + rs3
0x14	fdiv	Base	rd = rs1 / rs2
0x15	fneg	Base	rd = -rs1
0x16	fabs	Base	rd = abs(rs1)
0x17	fmin	Base	rd = min(rs1, rs2)
0x18	fmax	Base	rd = max(rs1, rs2)
0x19	fclamp	Extended	rd = clamp(rs1, rs2, rs3)
0x1A	fsqrt	Base	rd = sqrt(rs1)
0x1B	frsqrt	Base	rd = 1/sqrt(rs1)
0x1C	frcp	Base	rd = 1/rs1
0x1D	fround	Base	modifier: 0=floor, 1=ceil, 2=round, 3=trunc
0x1E	ffract	Base	rd = fract(rs1)
0x1F	ftransc	Base	modifier: 0=sin, 1=cos, 2=exp2, 3=log2

Bitwise operations: 0x20-0x27

Opcode	Mnemonic	Format	Description
0x20	and	Base	rd = rs1 & rs2
0x21	or	Base	rd = rs1 | rs2
0x22	xor	Base	rd = rs1 ^ rs2
0x23	not	Base	rd = ~rs1
0x24	shift	Base	modifier: 0=shl, 1=shr, 2=sar
0x25	bitop	Base	modifier: 0=bitcount, 1=bitfind, 2=bitrev
0x26	bfe	Extended	Extract bit field
0x27	bfi	Extended	Insert bit field

Comparison and select: 0x28-0x2F

Opcode	Mnemonic	Format	Description
0x28	icmp	Base	modifier: 0=eq, 1=ne, 2=lt, 3=le, 4=gt, 5=ge
0x29	ucmp	Base	modifier: 0=lt, 1=le
0x2A	fcmp	Base	modifier: 0=eq, 1=lt, 2=le, 3=gt, 4=ne, 5=ord, 6=unord
0x2B	select	Base	rd = pred ? rs1 : rs2
0x2C	fsat	Base	rd = clamp(rs1, 0.0, 1.0)
0x2D-0x2F	Reserved	:	:

Local memory: 0x30-0x37

Opcode	Mnemonic	Format	Description
0x30	local_load	Base	modifier: 0=u8, 1=u16, 2=u32, 3=u64
0x31	local_store	Base	modifier: 0=u8, 1=u16, 2=u32, 3=u64
0x32-0x37	Reserved	:	:

Device memory: 0x38-0x3F

Opcode	Mnemonic	Format	Description
0x38	device_load	Base	modifier: 0=u8, 1=u16, 2=u32, 3=u64, 4=u128
0x39	device_store	Base	modifier: 0=u8, 1=u16, 2=u32, 3=u64, 4=u128
0x3A-0x3F	Reserved	:	:

Atomic operations: 0x40-0x4F

Opcode	Mnemonic	Format	Description
0x40	atomic_add	Extended	Atomic add; scope in scope field
0x41	atomic_sub	Extended	Atomic subtract
0x42	atomic_min	Extended	Atomic minimum
0x43	atomic_max	Extended	Atomic maximum
0x44	atomic_and	Extended	Atomic bitwise AND
0x45	atomic_or	Extended	Atomic bitwise OR
0x46	atomic_xor	Extended	Atomic bitwise XOR
0x47	atomic_exchange	Extended	Atomic swap
0x48	atomic_cas	Extended	Compare-and-swap
0x49-0x4F	Reserved	:	:

Wave operations: 0x50-0x5F

Opcode	Mnemonic	Format	Description
0x50	wave_shuffle	Base	rd = rs1 from lane rs2
0x51	wave_shuffle_up	Base	rd = rs1 from lane (lane_id - rs2)
0x52	wave_shuffle_down	Base	rd = rs1 from lane (lane_id + rs2)
0x53	wave_shuffle_xor	Base	rd = rs1 from lane (lane_id ^ rs2)
0x54	wave_broadcast	Base	rd = rs1 from lane rs2; all threads
0x55	wave_ballot	Base	rd = bitmask of pd across active threads
0x56	wave_any	Base	pd_dst = any active thread has pd_src true
0x57	wave_all	Base	pd_dst = all active threads have pd_src true
0x58	wave_prefix_sum	Base	Exclusive prefix sum
0x59	wave_reduce	Base	modifier: 0=add, 1=min, 2=max
0x5A-0x5F	Reserved	:	:

Control flow and synchronization: 0x60-0x6F

Opcode	Mnemonic	Format	Description
0x60	if	Base	Begin conditional block
0x61	else	Base	Switch active/inactive
0x62	endif	Base	End conditional, restore mask
0x63	loop	Base	Begin loop
0x64	break	Base	Exit loop for predicated threads
0x65	continue	Base	Skip to next iteration for predicated threads
0x66	endloop	Base	End loop, branch back if any active
0x67	call	Extended	Call function at imm32 address
0x68	return	Base	Return from function
0x69	barrier	Base	Workgroup barrier
0x6A	fence	Base	modifier: 0=acquire, 1=release, 2=acq_rel
0x6B	wait	Base	Wait for async loads
0x6C	halt	Base	Terminate thread
0x6D-0x6F	Reserved	:	:

Type conversion: 0x70-0x7F

Opcode	Mnemonic	Format	Description
0x70	cvt_f32_i32	Base	Signed int to float
0x71	cvt_f32_u32	Base	Unsigned int to float
0x72	cvt_i32_f32	Base	Float to signed int
0x73	cvt_u32_f32	Base	Float to unsigned int
0x74	cvt_f32_f16	Base	F16 to F32
0x75	cvt_f16_f32	Base	F32 to F16
0x76	cvt_f32_f64	Base	F64 to F32; requires CAP_F64
0x77	cvt_f64_f32	Base	F32 to F64; requires CAP_F64
0x78-0x7F	Reserved	:	:

F16 arithmetic: 0x80-0x8F

Opcode	Mnemonic	Format	Description
0x80	hadd	Base	F16 add
0x81	hsub	Base	F16 subtract
0x82	hmul	Base	F16 multiply
0x83	hma	Extended	F16 fused multiply-add
0x84	hadd2	Base	Packed 2xF16 add
0x85	hmul2	Base	Packed 2xF16 multiply
0x86	hma2	Extended	Packed 2xF16 fused multiply-add
0x87-0x8F	Reserved	:	:

F64 arithmetic: 0x90-0x9F: requires CAP_F64

Opcode	Mnemonic	Format	Description
0x90	dadd	Base	F64 add
0x91	dsub	Base	F64 subtract
0x92	dmul	Base	F64 multiply
0x93	dma	Extended	F64 fused multiply-add
0x94	ddiv	Base	F64 divide
0x95	dsqrt	Base	F64 square root
0x96-0x9F	Reserved	:	:

Matrix MMA: 0xA0-0xAF: requires CAP_MMA

Opcode	Mnemonic	Format	Description
0xA0	mma_f16_f32	Extended	D = A*B+C, A/B F16, C/D F32
0xA1	mma_bf16_f32	Extended	D = A*B+C, A/B BF16, C/D F32
0xA2	mma_f32_f32	Extended	D = A*B+C, all F32
0xA3-0xAF	Reserved	:	:

Miscellaneous: 0xF0-0xFF

Opcode	Mnemonic	Format	Description
0xF0	mov	Base	rd = rs1
0xF1	mov_imm	Extended	rd = imm32
0xF2	mov_special	Base	rd = special register; rs1 encodes which
0xF3	nop	Base	No operation
0xF4-0xFF	Reserved	:	:

B. Vendor mapping reference

Abstract concept	NVIDIA	AMD RDNA	AMD CDNA	Intel	Apple
Core	SM	WGP	CU	Xe-core	GPU core
Wave	Warp; 32	Wavefront; 32	Wavefront; 64	Sub-group; 8-16	SIMD-group; 32
Register	PTX register	VGPR	VGPR	GRF entry	GPR
Local memory	Shared memory	LDS	LDS	SLM	Threadgroup mem
Barrier	bar.sync	S_BARRIER	S_BARRIER	barrier	threadgroup_barrier
Shuffle	__shfl_sync	DPP/ds_permute	DPP/ds_permute	sub_group_shuffle	simd_shuffle
Atomic	atom/red	ds/buffer atomic	ds/buffer atomic	atomic_ref; SEND	atomic_fetch_add
Device memory	Global memory	VRAM	VRAM	GDDR/HBM	LPDDR; unified
Fence	fence.scope	S_WAITCNT	S_WAITCNT	scoreboard	wait_for_loads