The WAVE Specification

Wide Architecture Virtual Encoding

Version: 0.2 (Working Draft, Revised)

Authors: Ojima Abraham, Onyinye Okoli

Date: March 23, 2026

Status: Working Draft (Revised)

0.Conformance Language

The key words MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL in this document are to be interpreted as described in RFC 2119.

1.Introduction

1.1Purpose

This specification defines WAVE (Wide Architecture Virtual Encoding), a vendor-neutral instruction set architecture for general-purpose GPU computation. It specifies an abstract execution model, a register model, a memory model, structured control flow semantics, an instruction set, and a capability query system.

The specification follows the thin abstraction principle: it defines what a compliant implementation MUST be able to do, not how it must do it. Implementations MAY use any microarchitectural technique to achieve compliance.

1.2Scope

This specification covers general-purpose compute workloads. Graphics pipeline operations (rasterization, tessellation, pixel export, ray tracing) are out of scope and MAY be addressed by future extensions.

1.3Design 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. Values that differ across implementations (wave width, register count, scratchpad size) are exposed as queryable constants, not fixed in the specification.
3. Structured divergence. The specification defines control flow semantics but not divergence mechanisms. Implementations are free to use any technique (execution masks, predication, hardware stacks, per-thread program counters) to achieve the specified behavior.
4. Mandatory minimums. Every queryable parameter has a minimum value. A compliant implementation MUST meet or exceed all minimums.

1.4Relationship to Other Standards

This specification is complementary to existing 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. The distinction is that this specification defines the hardware execution model, while existing standards define host-device interaction.

1.5Changes from v0.1

This version incorporates corrections and clarifications discovered during implementation of the reference toolchain (assembler, disassembler, and emulator). Key changes:

1. Register encoding widened from 5-bit to 8-bit (Section 8.2). The v0.1 encoding used 5-bit register fields, limiting addressable registers to 32. This conflicted with Section 7.1, which specifies MAX_REGISTERS minimum of 64. The encoding is now 8-bit, supporting up to 256 registers per thread.
2. Minimum divergence stack depth specified (Section 5.4, Section 7.1). v0.1 did not define a minimum nesting depth for structured control flow. Implementations MUST now support at least 32 levels of nested control flow.
3. Predicate negation semantics clarified (Section 5.1). v0.1 did not specify behavior of negated predicates on break and continue instructions. This is now explicitly defined.
4. Per-Wave control flow state requirement added (Section 5.5). v0.1 did not specify whether control flow state (divergence stack) is per-Wave or shared. Each Wave MUST have independent control flow state.
5. Full opcode table provided (Appendix A). v0.1 deferred this to a future version.
6. Conformance test suite referenced (Section 9.4). The reference test suite (102 tests) is now a companion artifact.

2.Execution Model

2.1Overview

A compliant processor consists of one or more Cores. Each Core is an independent compute unit capable of executing multiple Workgroups concurrently. Cores are not addressable by software. The hardware assigns Workgroups to Cores, and the programmer MUST NOT assume any particular mapping.

2.2Thread 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 a position within the hierarchy identified by hardware-populated identity values. 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 the purpose of 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 mechanism by which this is achieved is not specified.

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.

The number of Waves per Workgroup is ceil(workgroup_thread_count / W).

Workgroup dimensions are specified at dispatch time as a 3-dimensional size (x, y, 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.3Thread Identifiers

Every Thread has the following hardware-populated, read-only values available as special registers:

2.4Core Resources

Each Core provides the following resources:

Register File. A fixed-size on-chip storage of F bytes, partitioned among all simultaneously resident Threads. Each Thread receives R registers (declared at compile time). The maximum number of simultaneously resident Waves is bounded by:

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

where 4 is the register width in bytes (32 bits). This is the occupancy equation. Implementations MUST support at least MIN_MAX_REGISTERS registers per Thread.

Local Memory. A fixed-size on-chip scratchpad of S bytes, shared among all Waves in the same Workgroup. Local Memory is explicitly addressed via load and store instructions. There is no automatic caching or data placement. Local Memory contents are undefined at Workgroup start and are not preserved across Workgroup boundaries. Implementations MUST provide at least MIN_LOCAL_MEMORY_SIZE bytes.

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.5Execution 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. Period.
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 (no deadlock from scheduling).
7. A Wave that has been dispatched MUST eventually complete, assuming the program terminates (no starvation).
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.6Dispatch

A dispatch operation launches a Grid of Workgroups. The dispatch specifies:

• The kernel (program entry point)
• Grid dimensions
• Workgroup dimensions
• Register count per Thread (R)
• Local Memory size required (in bytes)
• Kernel arguments (buffer addresses, constants)

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

3.Register Model

3.1General-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 how the register contents are interpreted (integer, float, bitfield). Registers are named r0 through r{R-1}.

3.2Sub-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.3Register Pairs

Two consecutive registers MAY be used as a 64-bit value: r{N}:r{N+1} with r{N} as the low 32 bits. This is used for F64 operations (where supported) and 64-bit integer operations.

3.4Special 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.5Predicate 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.6Allocation and Occupancy

The register count R is declared per-kernel at compile time. The implementation allocates R registers per Thread for all Threads in all resident Waves. The occupancy equation determines how many Waves can be resident simultaneously. Compilers SHOULD minimize R to maximize occupancy.

4.Memory Model

4.1Memory 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.2Local 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. Accesses to the same address within a bank (broadcast) MUST NOT cause a conflict.

Supports access widths of 8, 16, 32, and 64 bits.

4.3Device 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.4Memory Ordering

Default ordering is relaxed. Ordering is achieved through scoped fence operations 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.5Atomic 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.1Structured 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.

Predicate negation on break and continue: When a break or continue instruction uses a negated predicate (e.g., break !p0), the instruction applies to Threads where the predicate is false. That is, break !p0 causes Threads where p0 is false to exit the loop. The negation is applied before evaluating which Threads are affected.

Function call:

call <function>
return

Function calls push the return address onto an implementation-managed call stack. The call stack depth MUST support at least MAX_CALL_DEPTH levels of nesting (see Section 7.1). Recursion is OPTIONAL (see Section 7).

5.2Uniform Branches

If all Threads in a Wave evaluate a branch identically (uniform), the implementation SHOULD avoid executing the not-taken path. Performance optimization, not correctness requirement.

5.3Divergence and Reconvergence

The implementation is free to use any mechanism to implement 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.4Minimum Divergence Depth

Implementations MUST support nested control flow (if/else/endif, loop/break/endloop) to a depth of at least MIN_DIVERGENCE_DEPTH levels (see Section 7.1). This means a program may have up to MIN_DIVERGENCE_DEPTH nested if/else/endif blocks, or nested loops, or any combination thereof.

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

5.5Per-Wave Control Flow State

Each Wave MUST maintain independent control flow state. This includes, but is not limited to, 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.

Rationale: This requirement was added in v0.2 after the reference emulator discovered that sharing control flow state across Waves in a Workgroup causes deadlock when Waves reach barriers at different loop iterations.

6.Instruction Set

6.1Instruction Format

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

Instructions are specified in this document in 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 where the predicate condition is met. Threads where the condition is not met are unaffected — their destination registers retain their previous values and no side effects (memory stores, atomics) occur.

The binary encoding of instructions is defined in Section 8.

6.2Integer Arithmetic

All integer operations are performed per-Thread. Full set of integer arithmetic for i32/u32:

• 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.3Bitwise 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.4Floating-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.5Floating-Point Arithmetic (F16) — REQUIRED

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

• hadd2, hmul2, hma2

6.6Floating-Point Arithmetic (F64) — OPTIONAL

F64 operations on register pairs:

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

6.7Type Conversion

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

6.8Comparison and Select

Comparison instructions set predicate registers. Select instruction for conditional moves.

6.9Memory 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.10Atomic Operations

Atomic instructions on Local and Device Memory with scope suffixes (.wave, .workgroup, .device, .system). Return old value; non-returning variants SHOULD be optimized.

6.11Wave 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

For shuffle operations, if the source lane is out of bounds (< 0 or >= W) or the source lane is inactive, the result is implementation-defined.

Wave operations operate only on active Threads. Inactive Threads (masked by divergence) do not participate in reductions, ballots, or prefix sums, and do not have their registers modified.

6.12Synchronization

• 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)

Barrier restriction: A barrier instruction MUST NOT appear inside a divergent control flow path. That is, when a Wave reaches a barrier, all active Threads in that Wave (at the point of the outermost non-divergent scope) must reach the same barrier. Barriers inside uniform if blocks (where all Threads agree) are permitted. Barriers inside loops are permitted provided all Waves in the Workgroup execute the same number of barrier instructions per loop iteration.

6.13Control Flow Instructions

6.14Matrix Multiply-Accumulate — OPTIONAL

• mma_f16_f32, mma_bf16_f32, mma_f32_f32

Tile dimensions queryable.

Miscellaneous

• mov, mov_imm, nop

7.Capability System

7.1Required Constants

7.2Optional 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.3Matrix 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.4Query Mechanism

Host API provides query_constant and query_capability functions.

8.Binary Encoding

8.1Overview

Instructions are encoded as fixed-width 48-bit (6-byte) words. Some instructions require an additional 32-bit word for immediate values or extended operands (80 bits / 10 bytes total).

v0.2 change: The v0.1 encoding used 32-bit base instructions with 5-bit register fields (max 32 registers). This conflicted with MAX_REGISTERS = 64. The encoding has been widened to 48 bits with 8-bit register fields (max 256 registers).

8.2Base Instruction Format (48-bit)

8.3Extended 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.4Opcode Map

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

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

9.Conformance

9.1Required 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.2Implementation-Defined Behavior

The following behaviors are implementation-defined (valid implementations may differ):

• 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.3Undefined Behavior

The following constitute undefined behavior (no guarantees):

• 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.4Conformance Testing

A reference conformance test suite consisting of 102 tests is provided as a companion artifact in the WAVE toolchain repository. The test suite verifies:

1. Correct execution of all mandatory instructions, including edge cases (overflow, NaN, 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 102 tests in the mandatory suite produce correct results. The test suite is versioned alongside the specification.

A.Full Opcode Table

Integer Arithmetic (0x00-0x0F)

Floating-Point Arithmetic F32 (0x10-0x1F)

Bitwise Operations (0x20-0x27)

Comparison and Select (0x28-0x2F)

Local Memory (0x30-0x37)

Device Memory (0x38-0x3F)

Atomic Operations (0x40-0x4F)

Wave Operations (0x50-0x5F)

Control Flow and Synchronization (0x60-0x6F)

Type Conversion (0x70-0x7F)

F16 Arithmetic (0x80-0x8F)

F64 Arithmetic (0x90-0x9F) — Requires CAP_F64

Matrix MMA (0xA0-0xAF) — Requires CAP_MMA

Miscellaneous (0xF0-0xFF)

B.Vendor Mapping Reference

C.Revision History

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

occupancy = min(
    max_waves_per_core,
    floor(register_file_size / (registers_per_thread * wave_width * 4)),
    floor(local_memory_size / local_memory_per_workgroup)
)

[predicate] opcode[.modifiers] destination, source1, source2, ...

mma.m16n8k8.f32.f16 {d0,d1,d2,d3}, {a0,a1}, {b0}, {c0,c1,c2,c3}

wave_result wave_get_capability(wave_device device, wave_capability cap, void* value, size_t size);

uint32_t wave_width;
wave_get_capability(device, WAVE_CAP_WAVE_WIDTH, &wave_width, sizeof(wave_width));

31      26 25    21 20    16 15    11 10     6 5       0
┌─────────┬────────┬────────┬────────┬────────┬─────────┐
│ opcode  │   rd   │   ra   │   rb   │  pred  │  flags  │
│ (6 bits)│(5 bits)│(5 bits)│(5 bits)│(5 bits)│ (6 bits)│
└─────────┴────────┴────────┴────────┴────────┴─────────┘

63      58 57    53 52    48 47    43 42    38 37    32
┌─────────┬────────┬────────┬────────┬────────┬────────┐
│  1 1 1  │ opcode │   rd   │   ra   │   rb   │   rc   │
│ (3 bits)│(5 bits)│(5 bits)│(5 bits)│(5 bits)│(5 bits)│
└─────────┴────────┴────────┴────────┴────────┴────────┘
31                                                     0
┌──────────────────────────────────────────────────────┐
│                     immediate                         │
│                     (32 bits)                         │
└──────────────────────────────────────────────────────┘