WAVE specification v0.1

• Version: 0.1, working draft
• Authors: Ojima Abraham, Onyinye Okoli
• Date: March 22, 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.

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.

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. Threads on the non-taken path of a divergent branch are inactive. They produce no architectural side effects: no memory writes or register updates visible to other threads.
3. All threads in a workgroup that reach a barrier will synchronize at that barrier before any proceeds past it.
4. Memory operations to local memory are visible to all threads in the workgroup after a barrier.
5. Memory operations to device memory are visible according to the memory ordering rules, Section 4.
6. Wave-level reductions, broadcasts, and shuffles operate only on active threads.
7. A compliant implementation MUST eventually make progress on at least one wave per core.

2.6 Dispatch

A dispatch specifies:

1. A compiled program containing instructions for this ISA
2. Grid dimensions: the number of workgroups in x, y, z
3. Workgroup dimensions: the number of threads in x, y, z
4. Kernel arguments: device memory pointers and constants

The implementation distributes workgroups to cores and creates waves within each workgroup. The distribution policy is implementation-defined.

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 be at least MIN_MAX_REGISTERS. Registers are 32 bits wide and are denoted r0 through r(R-1). Register contents are undefined at thread start unless explicitly initialized.

3.2 Sub-register access

A 32-bit register may be accessed as two 16-bit halves:

• r0.lo: bits [15:0]
• r0.hi: bits [31:16]

An optional capability provides access as four 8-bit bytes:

• r0.b0: bits [7:0]
• r0.b1: bits [15:8]
• r0.b2: bits [23:16]
• r0.b3: bits [31:24]

3.3 Register pairs

64-bit operations use register pairs. The notation r0:r1 denotes a 64-bit value where r0 holds the low 32 bits and r1 holds the high 32 bits. Register pairs MUST use adjacent even-odd register numbers, such as r0:r1, r2:r3, but not r1:r2.

3.4 Special registers

The following read-only registers are populated by hardware:

• %tid.x, %tid.y, %tid.z: thread ID within workgroup
• %wid: wave ID within workgroup
• %lid: lane ID within wave
• %ctaid.x, %ctaid.y, %ctaid.z: workgroup ID within grid
• %ntid.x, %ntid.y, %ntid.z: workgroup dimensions
• %nctaid.x, %nctaid.y, %nctaid.z: grid dimensions
• %nwaves: number of waves in this workgroup
• %clock: cycle counter, implementation-defined resolution

3.5 Predicate registers

Implementations MUST provide at least 8 predicate registers, denoted p0 through p7. Each predicate register holds a boolean value, 1 bit per thread. Predicate registers are used for conditional execution and branch conditions.

3.6 Allocation and occupancy

Compilers declare the number of registers required by a kernel. The runtime uses this to compute occupancy:

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

Higher occupancy generally improves latency hiding but reduces registers and local memory available per thread.

4. Memory model

4.1 Memory spaces

The specification defines four memory spaces:

Space	Scope	Lifetime	Typical implementation
Private	Single thread	Thread lifetime	Registers or stack
Local	Workgroup	Workgroup lifetime	On-chip SRAM
Device	All threads	Kernel lifetime; at least	VRAM
Constant	All threads	Kernel lifetime	Cached read-only

4.2 Local memory details

Local memory MUST support atomic operations. Access latency is implementation-defined but SHOULD be significantly lower than device memory. Bank conflicts MAY increase access latency but MUST NOT affect correctness. The implementation MUST provide at least MIN_LOCAL_MEMORY_SIZE bytes per workgroup.

4.3 Device memory details

Device memory is the largest but slowest memory space. It persists across kernel launches, until explicitly deallocated by the host. Device memory MUST support atomic operations. Coalescing combines multiple thread accesses into fewer memory transactions. It is implementation-defined but strongly encouraged.

4.4 Memory ordering

Within a single thread, memory operations appear to execute in program order. Across threads, the ordering is relaxed by default. Explicit ordering uses:

A fence instruction orders memory operations before and after it. Fences have a scope:

• fence.wave: ordering visible to all threads in the wave
• fence.workgroup: ordering visible to all threads in the workgroup
• fence.device: ordering visible to all threads on the device

The barrier instruction implies a workgroup-scope fence. All memory operations before the barrier are visible to all threads in the workgroup after the barrier.

4.5 Atomic operations

Atomic operations are indivisible read-modify-write operations. Implementations MUST support the following atomic operations on both local and device memory:

• atom.add, atom.sub: addition/subtraction
• atom.min, atom.max: minimum/maximum, signed and unsigned
• atom.and, atom.or, atom.xor: bitwise operations
• atom.exch: exchange
• atom.cas: compare-and-swap

64-bit atomics are an optional capability.

5. Control flow

5.1 Structured control flow

All control flow MUST be structured. The specification defines the following control flow primitives:

• if pd, label_else, label_endif: conditional branch based on predicate pd
• else: marks the start of the else block
• endif: marks the end of the if/else construct
• loop label_end: begins a loop, label_end marks the end
• endloop: marks the end of a loop
• break pd: exits the innermost loop if predicate is true
• continue pd: jumps to the next iteration of the innermost loop

Arbitrary goto is not supported. All control flow graphs MUST be reducible.

5.2 Uniform branches

A branch is uniform if all active threads in a wave take the same direction. Uniform branches have no divergence overhead. Compilers SHOULD annotate branches as uniform when known, allowing optimizations.

The syntax for annotating a branch as uniform is @uniform if pd, .... If a branch annotated as uniform diverges at runtime, behavior is undefined.

5.3 Divergence and reconvergence

When threads in a wave disagree on a branch condition, the wave is divergent. The implementation MUST execute both paths, but the order and mechanism are implementation-defined. Possible mechanisms include:

• Execution masks, execute both paths with some threads inactive
• Stack-based reconvergence
• Per-thread program counters

The specification only requires that:

1. All threads reconverge after structured control flow constructs, at endif, endloop.
2. Inactive threads produce no side effects.
3. Wave-level operations operate only on active threads.

6. Instruction set

6.1 Instruction format

All instructions follow the format:

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

Where:

• predicate, optional: @p0...@p7 for conditional execution
• opcode: instruction mnemonic
• modifiers: type suffixes, rounding modes, etc.
• destination: target register
• source: source operands, registers, immediates

6.2 Integer arithmetic

Instruction	Description	Operation
add.s32 rd, ra, rb	Signed 32-bit add	rd = ra + rb
add.u32 rd, ra, rb	Unsigned 32-bit add	rd = ra + rb
sub.s32 rd, ra, rb	Signed 32-bit subtract	rd = ra - rb
mul.lo.s32 rd, ra, rb	Signed multiply using the low 32 bits	rd = (ra * rb)[31:0]
mul.hi.s32 rd, ra, rb	Signed multiply using the high 32 bits	rd = (ra * rb)[63:32]
mul.wide.s32 rd:rd+1, ra, rb	Signed 32x32->64 multiply	rd:rd+1 = ra * rb
mad.lo.s32 rd, ra, rb, rc	Multiply-add using the low bits	rd = (ra * rb)[31:0] + rc
div.s32 rd, ra, rb	Signed division	rd = ra / rb
rem.s32 rd, ra, rb	Signed remainder	rd = ra % rb
neg.s32 rd, ra	Negate	rd = -ra
abs.s32 rd, ra	Absolute value	rd = |ra|
min.s32 rd, ra, rb	Minimum	rd = min(ra, rb)
max.s32 rd, ra, rb	Maximum	rd = max(ra, rb)

6.3 Bitwise operations

Instruction	Description	Operation
and.b32 rd, ra, rb	Bitwise AND	rd = ra & rb
or.b32 rd, ra, rb	Bitwise OR	rd = ra | rb
xor.b32 rd, ra, rb	Bitwise XOR	rd = ra ^ rb
not.b32 rd, ra	Bitwise NOT	rd = ~ra
shl.b32 rd, ra, rb	Shift left	rd = ra << rb
shr.u32 rd, ra, rb	Logical shift right	rd = ra >>> rb
shr.s32 rd, ra, rb	Arithmetic shift right	rd = ra >> rb
popc.b32 rd, ra	Population count	rd = popcount(ra)
clz.b32 rd, ra	Count leading zeros	rd = clz(ra)
brev.b32 rd, ra	Bit reverse	rd = reverse_bits(ra)
bfe.u32 rd, ra, rb, rc	Bit field extract	rd = (ra >> rb) & mask(rc)
bfi.b32 rd, ra, rb, rc, re	Bit field insert	Insert rc bits of ra at position rb into re

6.4 Floating-point: F32

All implementations MUST support IEEE 754-2008 single-precision, binary32, with the following exceptions:

• Denormal inputs MAY be flushed to zero.
• Denormal outputs MAY be flushed to zero.
• NaN payloads are implementation-defined.

Instruction	Description
add.f32 rd, ra, rb	Addition
sub.f32 rd, ra, rb	Subtraction
mul.f32 rd, ra, rb	Multiplication
fma.f32 rd, ra, rb, rc	Fused multiply-add: rd = ra*rb + rc
div.f32 rd, ra, rb	Division
rcp.f32 rd, ra	Reciprocal; approximate, 1 ULP
sqrt.f32 rd, ra	Square root
rsqrt.f32 rd, ra	Reciprocal square root; approximate
neg.f32 rd, ra	Negate
abs.f32 rd, ra	Absolute value
min.f32 rd, ra, rb	Minimum; IEEE semantics
max.f32 rd, ra, rb	Maximum; IEEE semantics
sin.f32 rd, ra	Sine; approximate
cos.f32 rd, ra	Cosine; approximate
exp2.f32 rd, ra	2^x; approximate
log2.f32 rd, ra	log2(x); approximate

6.5 Floating-point: F16

Half-precision support is an optional capability. When present:

• Two f16 values are packed into a single 32-bit register.
• Operations operate on pairs, vec2, or scalars.

Instruction	Description
add.f16x2 rd, ra, rb	Packed f16 addition
mul.f16x2 rd, ra, rb	Packed f16 multiplication
fma.f16x2 rd, ra, rb, rc	Packed f16 FMA

6.6 Floating-point: F64

Double-precision support is an optional capability. When present, operations use register pairs.

Instruction	Description
add.f64 rd:rd+1, ra:ra+1, rb:rb+1	f64 addition
mul.f64 rd:rd+1, ra:ra+1, rb:rb+1	f64 multiplication
fma.f64 rd:rd+1, ra:ra+1, rb:rb+1, rc:rc+1	f64 FMA
div.f64 rd:rd+1, ra:ra+1, rb:rb+1	f64 division

6.7 Type conversion

Instruction	Description
cvt.f32.s32 rd, ra	Signed int to float
cvt.f32.u32 rd, ra	Unsigned int to float
cvt.s32.f32 rd, ra	Float to signed int; truncate
cvt.rni.s32.f32 rd, ra	Float to signed int; round nearest
cvt.f16.f32 rd, ra	f32 to f16; packed
cvt.f32.f16 rd, ra	f16 to f32; unpacked
cvt.f64.f32 rd:rd+1, ra	f32 to f64
cvt.f32.f64 rd, ra:ra+1	f64 to f32

6.8 Comparison and select

Instruction	Description
setp.eq.s32 pd, ra, rb	Set predicate if equal
setp.ne.s32 pd, ra, rb	Set predicate if not equal
setp.lt.s32 pd, ra, rb	Set predicate if less than
setp.le.s32 pd, ra, rb	Set predicate if less or equal
setp.gt.s32 pd, ra, rb	Set predicate if greater than
setp.ge.s32 pd, ra, rb	Set predicate if greater or equal
selp.s32 rd, ra, rb, pd	Select: rd = pd ? ra: rb
slct.s32.f32 rd, ra, rb, rc	Select based on sign: rd = (rc >= 0) ? ra : rb

Floating-point comparisons include special handling for NaN:

• setp.eq.f32: false if either operand is NaN
• setp.neu.f32: true if either operand is NaN, unordered not-equal

6.9 Memory operations

Instruction	Description
ld.local.b32 rd, [ra]	Load 32 bits from local memory
ld.local.b64 rd:rd+1, [ra]	Load 64 bits from local memory
ld.global.b32 rd, [ra]	Load 32 bits from device memory
ld.const.b32 rd, [ra]	Load 32 bits from constant memory
st.local.b32 [ra], rb	Store 32 bits to local memory
st.global.b32 [ra], rb	Store 32 bits to device memory

Vector loads and stores:

• ld.global.v2.b32 {rd, rd+1}, [ra]: load 2x32 bits
• ld.global.v4.b32 {rd, rd+1, rd+2, rd+3}, [ra]: load 4x32 bits

6.10 Atomic operations

Instruction	Description
atom.local.add.u32 rd, [ra], rb	Atomic add to local memory
atom.global.add.u32 rd, [ra], rb	Atomic add to device memory
atom.global.cas.b32 rd, [ra], rb, rc	Compare-and-swap
atom.global.exch.b32 rd, [ra], rb	Exchange
atom.global.min.s32 rd, [ra], rb	Atomic minimum
atom.global.max.s32 rd, [ra], rb	Atomic maximum

All atomic operations return the previous value.

6.11 Wave operations

Wave operations compute across all active threads in a wave.

Instruction	Description
wave.reduce.add.u32 rd, ra	Sum of ra across all active lanes
wave.reduce.min.s32 rd, ra	Minimum of ra across all active lanes
wave.reduce.max.s32 rd, ra	Maximum of ra across all active lanes
wave.reduce.and.b32 rd, ra	Bitwise AND across all active lanes
wave.reduce.or.b32 rd, ra	Bitwise OR across all active lanes
wave.broadcast.b32 rd, ra, rb	Broadcast ra from lane rb to all lanes
wave.shuffle.b32 rd, ra, rb	rd = ra from lane rb
wave.shuffle.xor.b32 rd, ra, rb	rd = ra from lane (lane_id ^ rb)
wave.shuffle.up.b32 rd, ra, rb	rd = ra from lane (lane_id - rb)
wave.shuffle.down.b32 rd, ra, rb	rd = ra from lane (lane_id + rb)
wave.prefix.add.u32 rd, ra	Exclusive prefix sum
wave.ballot.b32 rd, pd	rd = bitmask of pd across all lanes
wave.any pd, ps	pd = true if any active lane has ps=true
wave.all pd, ps	pd = true if all active lanes have ps=true

6.12 Synchronization

Instruction	Description
barrier	Workgroup barrier + memory fence
fence.wave	Memory fence, wave scope
fence.workgroup	Memory fence, workgroup scope
fence.device	Memory fence, Device scope

6.13 Control flow instructions

Instruction	Description
if pd	Begin if block
else	Begin else block
endif	End if/else block
loop	Begin loop
endloop	End loop
break pd	Break from loop if predicate true
continue pd	Continue to next iteration if predicate true
ret	Return from kernel

6.14 Matrix MMA: optional

Matrix multiply-accumulate, MMA, operations are an optional capability. When present, implementations support tensor core-style operations:

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

Shapes and data types are queryable through the capability system. See Section 7.

7. Capability system

7.1 Required constants

Every implementation MUST provide the following queryable constants:

Constant	Minimum	Description
WAVE_WIDTH	16	Number of threads per wave
MAX_WORKGROUP_SIZE	256	Maximum threads per workgroup
MAX_REGISTERS	64	Maximum registers per thread
LOCAL_MEMORY_SIZE	16384	Bytes of local memory per workgroup
MAX_WAVES_PER_CORE	16	Maximum resident waves per core
PREDICATE_REGISTERS	8	Number of predicate registers
CLUSTER_SIZE	1	Workgroups per cluster; 1 = no clusters

7.2 Optional capabilities

Implementations MAY support the following optional features:

Capability	Description
CAP_F16	Half-precision floating-point
CAP_F64	Double-precision floating-point
CAP_ATOMIC64	64-bit atomic operations
CAP_MMA	Matrix multiply-accumulate
CAP_DP4A	4-element dot product; int8
CAP_SUBGROUPS	Subgroup operations; partial wave
CAP_CLUSTER	Cluster-scope operations

7.3 MMA parameters

If CAP_MMA is present, the following constants define supported MMA shapes:

Constant	Description
MMA_M	M dimension; rows of output
MMA_N	N dimension; columns of output
MMA_K	K dimension; inner dimension
MMA_INPUT_TYPES	Bitmask of supported input types
MMA_OUTPUT_TYPES	Bitmask of supported output types

7.4 Query mechanism

The host runtime provides a query function:

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

The cap argument enumerates all queryable constants and capabilities. The function writes the value to the provided buffer and returns a success/error code.

Example:

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

8. Binary encoding

8.1 Overview

Instructions are encoded in either a 32-bit base format or a 64-bit extended format. The base format accommodates common operations with register operands. The extended format adds support for immediate values, additional operands, and modifiers.

8.2 Base format: 32-bit

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

• opcode: Primary operation code, 64 base opcodes
• rd: Destination register
• ra, rb: Source registers
• pred: Predicate register, 0 = unconditional, 1-7 = p1-p7, 8-15 = !p0-!p7
• flags: Operation-specific flags, saturation, rounding, etc.

8.3 Extended format: 64-bit

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

Opcode bits [31:29] = 111 indicate the extended format, which supports:

• 4-operand instructions, FMA, MMA
• 32-bit immediate values
• Extended opcodes, 224 additional operations

9. Conformance

9.1 Required behavior

A conformant implementation MUST:

1. Execute all required instructions with correct semantics.
2. Meet or exceed all minimum capability values.
3. Provide correct memory ordering per Section 4.
4. Correctly handle divergence per Section 5.
5. Report capabilities accurately via Section 7.

9.2 Implementation-defined

The following are implementation-defined:

• Scheduling policy, wave selection, workgroup assignment
• Memory coalescing behavior
• Bank conflict penalties
• Divergence mechanism
• NaN payload values
• Clock counter resolution

9.3 Undefined behavior

The following have undefined behavior:

• Barrier within divergent control flow
• Uniform branch that diverges at runtime
• Access to unallocated registers
• Out-of-bounds memory access
• Data races: concurrent access without synchronization where at least one access is a write

9.4 Conformance testing

The planned reference test suite will include:

• Instruction correctness tests for all required operations
• Memory ordering tests
• Divergence correctness tests
• Capability query tests
• Edge case tests, overflow, NaN handling, etc.

A. Full opcode table

Deferred to a future version of this specification.

B. Vendor mapping

This appendix provides suggested mappings to vendor ISAs for reference implementations:

WAVE concept	NVIDIA; PTX	AMD; RDNA	Intel; Xe	Apple; M-series
Wave	Warp; 32	Wave; 32/64	SIMD; 8/16	SIMD-group; 32
Workgroup	Thread block	Work-group	Thread group	Threadgroup
Local memory	Shared memory	LDS	SLM	Threadgroup memory
Device memory	Global memory	VRAM	Global memory	Device memory
wave.shuffle	shfl.sync	ds_permute	mov; cross-lane	quad_shuffle
Barrier	bar.sync	s_barrier	barrier	threadgroup_barrier
Fence	fence.sc	S_WAITCNT	scoreboard	wait_for_loads