Binary encoding

Each WAVE instruction contains one or two 32-bit words.

Word0: 32-bit base format

Every WAVE instruction starts with a 32-bit word with the following bit layout:

Bits	Field	Width	Purpose
31:24	opcode	8 bits	Operation class
23:16	RD	8 bits	Destination register; 0 through 255
15:8	RS1	8 bits	Source register 1; 0 through 255
7:4	modifier	4 bits	Operation variant within opcode class; 0-15
3	reserved	1 bit	Reserved, must be zero
2	pred_neg	1 bit	Negate predicate condition
1:0	pred_reg	2 bits	Predicate selector; zero without negation means unpredicated execution

 31    24 23    16 15     8 7    4  3  2  1  0
┌────────┬────────┬────────┬──────┬───┬───┬────┐
│ opcode │   RD   │  RS1   │ mod  │rsv│neg│pred│
│  8b    │  8b    │  8b    │ 4b   │1b │1b │ 2b │
└────────┴────────┴────────┴──────┴───┴───┴────┘

Predicate encoding

When pred_reg=0 and pred_neg=0, the instruction is unpredicated. A nonzero selector tests the named predicate register. pred_neg=1 negates that test. For example, @p1 sets pred_reg=1 and pred_neg=0; @!p2 sets pred_reg=2 and pred_neg=1.

Predicate register p0 is writable and can be used by comparisons and control-flow instructions. See Encoding and execution invariants.

Register fields

Eight bits address 256 registers. Word0 holds rd and rs1; word1 holds additional source registers.

Modifier field

The 4-bit modifier selects one of up to 16 variants within an opcode. See Modifier field for the assignments.

Opcode 0x3F uses modifiers 0 through 7 for if, else, endif, loop, break, continue, endloop, and call. Modifiers 8 through 15 select return, halt, barrier, the fence variants, wait, and nop. Synchronization modifiers use SYNC_MODIFIER_OFFSET = 8.

Word1: 32-bit extended format

When an instruction needs additional source registers, a memory scope, or an inline immediate, it uses a second 32-bit word. The opcode in word0 determines whether word1 is present.

Bits	Field	Width	Purpose
31:24	RS2	8 bits	Source register 2; 0 through 255
23:16	RS3	8 bits	Source register 3; 0 through 255
15:8	RS4	8 bits	Source register 4; 0 through 255
7:2	reserved	6 bits	Reserved, must be zero
1:0	scope	2 bits	Memory scope; 00=wave, 01=workgroup, 10=device, 11=system

 31    24 23    16 15     8 7         2 1  0
┌────────┬────────┬────────┬──────────┬────┐
│  RS2   │  RS3   │  RS4   │ reserved │scop│
│  8b    │  8b    │  8b    │   6b     │ 2b │
└────────┴────────┴────────┴──────────┴────┘

Alternatively, the entire word1 can serve as a 32-bit immediate value, such as for mov_imm. The opcode determines interpretation.

Scope encoding

Word1 bits [1:0] select the scope of device atomics and fences:

Value	Scope	Meaning
00	Wave	Visible within the executing wave
01	Workgroup	Visible to all waves in the workgroup
10	Device	Visible to all waves on the device
11	System	Visible to the device and host CPU

When to use 64-bit encoding

• Two-source operations such as iadd r0, r1, r2 use RS1 in word0 and RS2 in word1.
• Fused multiply-add, fma r0, r1, r2, r3: needs RS1, RS2, and RS3.
• Immediate loads, mov_imm r0, 0x3F800000: the 32-bit immediate occupies the entire word1.
• Scoped atomics, atomic_add r0, r1, r2, workgroup: scope in word1 bits [1:0].
• Atomic compare-and-swap: needs address, expected, desired, and result registers.

WBIN container format

The .wbin container stores compiled WAVE programs. The parser validates section bounds and kernel ranges before exposing code.

Header: 32 bytes

Offset	Size	Field	Description
0	4	magic	ASCII "WAVE"; 0x57415645
4	2	version	Container format version
6	2	flags	Container flags
8	4	code_offset	Byte offset to code section
12	4	code_size	Size of code section in bytes
16	4	symbol_offset	Byte offset to symbol table
20	4	symbol_size	Size of symbol table in bytes
24	4	metadata_offset	Byte offset to metadata section
28	4	metadata_size	Size of metadata section in bytes

Kernel metadata

Each kernel entry in the metadata section describes the resource requirements that a backend needs for dispatch:

Field	Purpose
name	Kernel identifier; null-terminated UTF-8
register_count	Number of registers used per thread
local_mem_size	Bytes of workgroup-local memory required
workgroup_size_x	Workgroup dimension X
workgroup_size_y	Workgroup dimension Y
workgroup_size_z	Workgroup dimension Z
code_offset	Byte offset of this kernel's code within the code section
code_size	Size of this kernel's code in bytes

The metadata section starts with a 32-bit kernel count, followed by 32-byte entries. Each entry stores the kernel name's offset into the symbol table and the fields listed above.