wave_decode/
opcodes.rs

1// Copyright 2026 Ojima Abraham
2// SPDX-License-Identifier: Apache-2.0
3
4//! Opcode definitions and constants for WAVE ISA decoding.
5//!
6//! Mirrors the exact opcode assignments from wave-asm to ensure
7//! consistency between assembler, disassembler, and emulator.
8
9pub const OPCODE_SHIFT: u32 = 26;
10pub const OPCODE_MASK: u32 = 0x3F;
11pub const RD_SHIFT: u32 = 21;
12pub const RD_MASK: u32 = 0x1F;
13pub const RS1_SHIFT: u32 = 16;
14pub const RS1_MASK: u32 = 0x1F;
15pub const RS2_SHIFT: u32 = 11;
16pub const RS2_MASK: u32 = 0x1F;
17pub const MODIFIER_SHIFT: u32 = 7;
18pub const MODIFIER_MASK: u32 = 0x0F;
19pub const SCOPE_SHIFT: u32 = 5;
20pub const SCOPE_MASK: u32 = 0x03;
21pub const PRED_SHIFT: u32 = 3;
22pub const PRED_MASK: u32 = 0x03;
23pub const PRED_NEG_SHIFT: u32 = 2;
24pub const PRED_NEG_MASK: u32 = 0x01;
25pub const FLAGS_SHIFT: u32 = 0;
26pub const FLAGS_MASK: u32 = 0x03;
27pub const EXTENDED_RS3_SHIFT: u32 = 27;
28pub const EXTENDED_RS3_MASK: u32 = 0x1F;
29pub const EXTENDED_RS4_SHIFT: u32 = 22;
30pub const EXTENDED_RS4_MASK: u32 = 0x1F;
31pub const SYNC_OP_FLAG: u8 = 0x01;
32pub const MISC_OP_FLAG: u8 = 0x02;
33#[derive(Debug, Clone, Copy, PartialEq, Eq)]
34#[repr(u8)]
35pub enum Opcode {
36    Iadd = 0x00,
37    Isub = 0x01,
38    Imul = 0x02,
39    ImulHi = 0x03,
40    Imad = 0x04,
41    Idiv = 0x05,
42    Imod = 0x06,
43    Ineg = 0x07,
44    Iabs = 0x08,
45    Imin = 0x09,
46    Imax = 0x0A,
47    Iclamp = 0x0B,
48    Fadd = 0x10,
49    Fsub = 0x11,
50    Fmul = 0x12,
51    Fma = 0x13,
52    Fdiv = 0x14,
53    Fneg = 0x15,
54    Fabs = 0x16,
55    Fmin = 0x17,
56    Fmax = 0x18,
57    Fclamp = 0x19,
58    Fsqrt = 0x1A,
59    FUnaryOps = 0x1B,
60    F16Ops = 0x1C,
61    F16PackedOps = 0x1D,
62    F64Ops = 0x1E,
63    F64DivSqrt = 0x1F,
64    And = 0x20,
65    Or = 0x21,
66    Xor = 0x22,
67    Not = 0x23,
68    Shl = 0x24,
69    Shr = 0x25,
70    Sar = 0x26,
71    BitOps = 0x27,
72    Icmp = 0x28,
73    Ucmp = 0x29,
74    Fcmp = 0x2A,
75    Select = 0x2B,
76    Cvt = 0x2C,
77    LocalLoad = 0x30,
78    LocalStore = 0x31,
79    DeviceLoad = 0x38,
80    DeviceStore = 0x39,
81    LocalAtomic = 0x3C,
82    DeviceAtomic = 0x3D,
83    WaveOp = 0x3E,
84    Control = 0x3F,
85}
86impl Opcode {
87    #[must_use]
88    pub fn from_u8(value: u8) -> Option<Self> {
89        match value {
90            0x00 => Some(Self::Iadd),
91            0x01 => Some(Self::Isub),
92            0x02 => Some(Self::Imul),
93            0x03 => Some(Self::ImulHi),
94            0x04 => Some(Self::Imad),
95            0x05 => Some(Self::Idiv),
96            0x06 => Some(Self::Imod),
97            0x07 => Some(Self::Ineg),
98            0x08 => Some(Self::Iabs),
99            0x09 => Some(Self::Imin),
100            0x0A => Some(Self::Imax),
101            0x0B => Some(Self::Iclamp),
102            0x10 => Some(Self::Fadd),
103            0x11 => Some(Self::Fsub),
104            0x12 => Some(Self::Fmul),
105            0x13 => Some(Self::Fma),
106            0x14 => Some(Self::Fdiv),
107            0x15 => Some(Self::Fneg),
108            0x16 => Some(Self::Fabs),
109            0x17 => Some(Self::Fmin),
110            0x18 => Some(Self::Fmax),
111            0x19 => Some(Self::Fclamp),
112            0x1A => Some(Self::Fsqrt),
113            0x1B => Some(Self::FUnaryOps),
114            0x1C => Some(Self::F16Ops),
115            0x1D => Some(Self::F16PackedOps),
116            0x1E => Some(Self::F64Ops),
117            0x1F => Some(Self::F64DivSqrt),
118            0x20 => Some(Self::And),
119            0x21 => Some(Self::Or),
120            0x22 => Some(Self::Xor),
121            0x23 => Some(Self::Not),
122            0x24 => Some(Self::Shl),
123            0x25 => Some(Self::Shr),
124            0x26 => Some(Self::Sar),
125            0x27 => Some(Self::BitOps),
126            0x28 => Some(Self::Icmp),
127            0x29 => Some(Self::Ucmp),
128            0x2A => Some(Self::Fcmp),
129            0x2B => Some(Self::Select),
130            0x2C => Some(Self::Cvt),
131            0x30 => Some(Self::LocalLoad),
132            0x31 => Some(Self::LocalStore),
133            0x38 => Some(Self::DeviceLoad),
134            0x39 => Some(Self::DeviceStore),
135            0x3C => Some(Self::LocalAtomic),
136            0x3D => Some(Self::DeviceAtomic),
137            0x3E => Some(Self::WaveOp),
138            0x3F => Some(Self::Control),
139            _ => None,
140        }
141    }
142    #[must_use]
143    pub fn is_extended(self) -> bool {
144        matches!(
145            self,
146            Self::Imad
147                | Self::Iclamp
148                | Self::Fma
149                | Self::Fclamp
150                | Self::F16Ops
151                | Self::F16PackedOps
152                | Self::F64Ops
153                | Self::BitOps
154                | Self::LocalAtomic
155                | Self::DeviceAtomic
156                | Self::Control
157        )
158    }
159}
160#[derive(Debug, Clone, Copy, PartialEq, Eq)]
161#[repr(u8)]
162pub enum FUnaryOp {
163    Frsqrt = 0,
164    Frcp = 1,
165    Ffloor = 2,
166    Fceil = 3,
167    Fround = 4,
168    Ftrunc = 5,
169    Ffract = 6,
170    Fsat = 7,
171    Fsin = 8,
172    Fcos = 9,
173    Fexp2 = 10,
174    Flog2 = 11,
175}
176impl FUnaryOp {
177    #[must_use]
178    pub fn from_u8(value: u8) -> Option<Self> {
179        match value {
180            0 => Some(Self::Frsqrt),
181            1 => Some(Self::Frcp),
182            2 => Some(Self::Ffloor),
183            3 => Some(Self::Fceil),
184            4 => Some(Self::Fround),
185            5 => Some(Self::Ftrunc),
186            6 => Some(Self::Ffract),
187            7 => Some(Self::Fsat),
188            8 => Some(Self::Fsin),
189            9 => Some(Self::Fcos),
190            10 => Some(Self::Fexp2),
191            11 => Some(Self::Flog2),
192            _ => None,
193        }
194    }
195    #[must_use]
196    pub fn mnemonic(self) -> &'static str {
197        match self {
198            Self::Frsqrt => "frsqrt",
199            Self::Frcp => "frcp",
200            Self::Ffloor => "ffloor",
201            Self::Fceil => "fceil",
202            Self::Fround => "fround",
203            Self::Ftrunc => "ftrunc",
204            Self::Ffract => "ffract",
205            Self::Fsat => "fsat",
206            Self::Fsin => "fsin",
207            Self::Fcos => "fcos",
208            Self::Fexp2 => "fexp2",
209            Self::Flog2 => "flog2",
210        }
211    }
212}
213#[derive(Debug, Clone, Copy, PartialEq, Eq)]
214#[repr(u8)]
215pub enum F16Op {
216    Hadd = 0,
217    Hsub = 1,
218    Hmul = 2,
219    Hma = 3,
220}
221impl F16Op {
222    #[must_use]
223    pub fn from_u8(value: u8) -> Option<Self> {
224        match value {
225            0 => Some(Self::Hadd),
226            1 => Some(Self::Hsub),
227            2 => Some(Self::Hmul),
228            3 => Some(Self::Hma),
229            _ => None,
230        }
231    }
232    #[must_use]
233    pub fn mnemonic(self) -> &'static str {
234        match self {
235            Self::Hadd => "hadd",
236            Self::Hsub => "hsub",
237            Self::Hmul => "hmul",
238            Self::Hma => "hma",
239        }
240    }
241}
242#[derive(Debug, Clone, Copy, PartialEq, Eq)]
243#[repr(u8)]
244pub enum F16PackedOp {
245    Hadd2 = 0,
246    Hmul2 = 1,
247    Hma2 = 2,
248}
249impl F16PackedOp {
250    #[must_use]
251    pub fn from_u8(value: u8) -> Option<Self> {
252        match value {
253            0 => Some(Self::Hadd2),
254            1 => Some(Self::Hmul2),
255            2 => Some(Self::Hma2),
256            _ => None,
257        }
258    }
259    #[must_use]
260    pub fn mnemonic(self) -> &'static str {
261        match self {
262            Self::Hadd2 => "hadd2",
263            Self::Hmul2 => "hmul2",
264            Self::Hma2 => "hma2",
265        }
266    }
267}
268#[derive(Debug, Clone, Copy, PartialEq, Eq)]
269#[repr(u8)]
270pub enum F64Op {
271    Dadd = 0,
272    Dsub = 1,
273    Dmul = 2,
274    Dma = 3,
275}
276impl F64Op {
277    #[must_use]
278    pub fn from_u8(value: u8) -> Option<Self> {
279        match value {
280            0 => Some(Self::Dadd),
281            1 => Some(Self::Dsub),
282            2 => Some(Self::Dmul),
283            3 => Some(Self::Dma),
284            _ => None,
285        }
286    }
287    #[must_use]
288    pub fn mnemonic(self) -> &'static str {
289        match self {
290            Self::Dadd => "dadd",
291            Self::Dsub => "dsub",
292            Self::Dmul => "dmul",
293            Self::Dma => "dma",
294        }
295    }
296}
297#[derive(Debug, Clone, Copy, PartialEq, Eq)]
298#[repr(u8)]
299pub enum F64DivSqrtOp {
300    Ddiv = 0,
301    Dsqrt = 1,
302}
303impl F64DivSqrtOp {
304    #[must_use]
305    pub fn from_u8(value: u8) -> Option<Self> {
306        match value {
307            0 => Some(Self::Ddiv),
308            1 => Some(Self::Dsqrt),
309            _ => None,
310        }
311    }
312    #[must_use]
313    pub fn mnemonic(self) -> &'static str {
314        match self {
315            Self::Ddiv => "ddiv",
316            Self::Dsqrt => "dsqrt",
317        }
318    }
319}
320#[derive(Debug, Clone, Copy, PartialEq, Eq)]
321#[repr(u8)]
322pub enum BitOpType {
323    Bitcount = 0,
324    Bitfind = 1,
325    Bitrev = 2,
326    Bfe = 3,
327    Bfi = 4,
328}
329impl BitOpType {
330    #[must_use]
331    pub fn from_u8(value: u8) -> Option<Self> {
332        match value {
333            0 => Some(Self::Bitcount),
334            1 => Some(Self::Bitfind),
335            2 => Some(Self::Bitrev),
336            3 => Some(Self::Bfe),
337            4 => Some(Self::Bfi),
338            _ => None,
339        }
340    }
341    #[must_use]
342    pub fn mnemonic(self) -> &'static str {
343        match self {
344            Self::Bitcount => "bitcount",
345            Self::Bitfind => "bitfind",
346            Self::Bitrev => "bitrev",
347            Self::Bfe => "bfe",
348            Self::Bfi => "bfi",
349        }
350    }
351}
352#[derive(Debug, Clone, Copy, PartialEq, Eq)]
353#[repr(u8)]
354pub enum CmpOp {
355    Eq = 0,
356    Ne = 1,
357    Lt = 2,
358    Le = 3,
359    Gt = 4,
360    Ge = 5,
361    Ord = 6,
362    Unord = 7,
363}
364impl CmpOp {
365    #[must_use]
366    pub fn from_u8(value: u8) -> Option<Self> {
367        match value {
368            0 => Some(Self::Eq),
369            1 => Some(Self::Ne),
370            2 => Some(Self::Lt),
371            3 => Some(Self::Le),
372            4 => Some(Self::Gt),
373            5 => Some(Self::Ge),
374            6 => Some(Self::Ord),
375            7 => Some(Self::Unord),
376            _ => None,
377        }
378    }
379    #[must_use]
380    pub fn suffix(self) -> &'static str {
381        match self {
382            Self::Eq => "eq",
383            Self::Ne => "ne",
384            Self::Lt => "lt",
385            Self::Le => "le",
386            Self::Gt => "gt",
387            Self::Ge => "ge",
388            Self::Ord => "ord",
389            Self::Unord => "unord",
390        }
391    }
392}
393#[derive(Debug, Clone, Copy, PartialEq, Eq)]
394#[repr(u8)]
395pub enum CvtType {
396    F32I32 = 0,
397    F32U32 = 1,
398    I32F32 = 2,
399    U32F32 = 3,
400    F32F16 = 4,
401    F16F32 = 5,
402    F32F64 = 6,
403    F64F32 = 7,
404}
405impl CvtType {
406    #[must_use]
407    pub fn from_u8(value: u8) -> Option<Self> {
408        match value {
409            0 => Some(Self::F32I32),
410            1 => Some(Self::F32U32),
411            2 => Some(Self::I32F32),
412            3 => Some(Self::U32F32),
413            4 => Some(Self::F32F16),
414            5 => Some(Self::F16F32),
415            6 => Some(Self::F32F64),
416            7 => Some(Self::F64F32),
417            _ => None,
418        }
419    }
420    #[must_use]
421    pub fn mnemonic(self) -> &'static str {
422        match self {
423            Self::F32I32 => "cvt_f32_i32",
424            Self::F32U32 => "cvt_f32_u32",
425            Self::I32F32 => "cvt_i32_f32",
426            Self::U32F32 => "cvt_u32_f32",
427            Self::F32F16 => "cvt_f32_f16",
428            Self::F16F32 => "cvt_f16_f32",
429            Self::F32F64 => "cvt_f32_f64",
430            Self::F64F32 => "cvt_f64_f32",
431        }
432    }
433}
434#[derive(Debug, Clone, Copy, PartialEq, Eq)]
435#[repr(u8)]
436pub enum MemWidth {
437    U8 = 0,
438    U16 = 1,
439    U32 = 2,
440    U64 = 3,
441    U128 = 4,
442}
443impl MemWidth {
444    #[must_use]
445    pub fn from_u8(value: u8) -> Option<Self> {
446        match value {
447            0 => Some(Self::U8),
448            1 => Some(Self::U16),
449            2 => Some(Self::U32),
450            3 => Some(Self::U64),
451            4 => Some(Self::U128),
452            _ => None,
453        }
454    }
455    #[must_use]
456    pub fn suffix(self) -> &'static str {
457        match self {
458            Self::U8 => "u8",
459            Self::U16 => "u16",
460            Self::U32 => "u32",
461            Self::U64 => "u64",
462            Self::U128 => "u128",
463        }
464    }
465}
466#[derive(Debug, Clone, Copy, PartialEq, Eq)]
467#[repr(u8)]
468pub enum AtomicOp {
469    Add = 0,
470    Sub = 1,
471    Min = 2,
472    Max = 3,
473    And = 4,
474    Or = 5,
475    Xor = 6,
476    Exchange = 7,
477}
478impl AtomicOp {
479    #[must_use]
480    pub fn from_u8(value: u8) -> Option<Self> {
481        match value {
482            0 => Some(Self::Add),
483            1 => Some(Self::Sub),
484            2 => Some(Self::Min),
485            3 => Some(Self::Max),
486            4 => Some(Self::And),
487            5 => Some(Self::Or),
488            6 => Some(Self::Xor),
489            7 => Some(Self::Exchange),
490            _ => None,
491        }
492    }
493    #[must_use]
494    pub fn suffix(self) -> &'static str {
495        match self {
496            Self::Add => "add",
497            Self::Sub => "sub",
498            Self::Min => "min",
499            Self::Max => "max",
500            Self::And => "and",
501            Self::Or => "or",
502            Self::Xor => "xor",
503            Self::Exchange => "exchange",
504        }
505    }
506}
507#[derive(Debug, Clone, Copy, PartialEq, Eq)]
508#[repr(u8)]
509pub enum Scope {
510    Wave = 0,
511    Workgroup = 1,
512    Device = 2,
513    System = 3,
514}
515impl Scope {
516    #[must_use]
517    pub fn from_u8(value: u8) -> Option<Self> {
518        match value {
519            0 => Some(Self::Wave),
520            1 => Some(Self::Workgroup),
521            2 => Some(Self::Device),
522            3 => Some(Self::System),
523            _ => None,
524        }
525    }
526    #[must_use]
527    pub fn name(self) -> &'static str {
528        match self {
529            Self::Wave => "wave",
530            Self::Workgroup => "workgroup",
531            Self::Device => "device",
532            Self::System => "system",
533        }
534    }
535}
536#[derive(Debug, Clone, Copy, PartialEq, Eq)]
537#[repr(u8)]
538pub enum WaveOpType {
539    Shuffle = 0,
540    ShuffleUp = 1,
541    ShuffleDown = 2,
542    ShuffleXor = 3,
543    Broadcast = 4,
544    Ballot = 5,
545    Any = 6,
546    All = 7,
547}
548impl WaveOpType {
549    #[must_use]
550    pub fn from_u8(value: u8) -> Option<Self> {
551        match value {
552            0 => Some(Self::Shuffle),
553            1 => Some(Self::ShuffleUp),
554            2 => Some(Self::ShuffleDown),
555            3 => Some(Self::ShuffleXor),
556            4 => Some(Self::Broadcast),
557            5 => Some(Self::Ballot),
558            6 => Some(Self::Any),
559            7 => Some(Self::All),
560            _ => None,
561        }
562    }
563    #[must_use]
564    pub fn mnemonic(self) -> &'static str {
565        match self {
566            Self::Shuffle => "wave_shuffle",
567            Self::ShuffleUp => "wave_shuffle_up",
568            Self::ShuffleDown => "wave_shuffle_down",
569            Self::ShuffleXor => "wave_shuffle_xor",
570            Self::Broadcast => "wave_broadcast",
571            Self::Ballot => "wave_ballot",
572            Self::Any => "wave_any",
573            Self::All => "wave_all",
574        }
575    }
576}
577#[derive(Debug, Clone, Copy, PartialEq, Eq)]
578#[repr(u8)]
579pub enum WaveReduceType {
580    PrefixSum = 0,
581    ReduceAdd = 1,
582    ReduceMin = 2,
583    ReduceMax = 3,
584}
585impl WaveReduceType {
586    #[must_use]
587    pub fn from_u8(value: u8) -> Option<Self> {
588        match value {
589            0 => Some(Self::PrefixSum),
590            1 => Some(Self::ReduceAdd),
591            2 => Some(Self::ReduceMin),
592            3 => Some(Self::ReduceMax),
593            _ => None,
594        }
595    }
596    #[must_use]
597    pub fn mnemonic(self) -> &'static str {
598        match self {
599            Self::PrefixSum => "wave_prefix_sum",
600            Self::ReduceAdd => "wave_reduce_add",
601            Self::ReduceMin => "wave_reduce_min",
602            Self::ReduceMax => "wave_reduce_max",
603        }
604    }
605}
606#[derive(Debug, Clone, Copy, PartialEq, Eq)]
607#[repr(u8)]
608pub enum ControlOp {
609    If = 0,
610    Else = 1,
611    Endif = 2,
612    Loop = 3,
613    Break = 4,
614    Continue = 5,
615    Endloop = 6,
616    Call = 7,
617}
618impl ControlOp {
619    #[must_use]
620    pub fn from_u8(value: u8) -> Option<Self> {
621        match value {
622            0 => Some(Self::If),
623            1 => Some(Self::Else),
624            2 => Some(Self::Endif),
625            3 => Some(Self::Loop),
626            4 => Some(Self::Break),
627            5 => Some(Self::Continue),
628            6 => Some(Self::Endloop),
629            7 => Some(Self::Call),
630            _ => None,
631        }
632    }
633    #[must_use]
634    pub fn mnemonic(self) -> &'static str {
635        match self {
636            Self::If => "if",
637            Self::Else => "else",
638            Self::Endif => "endif",
639            Self::Loop => "loop",
640            Self::Break => "break",
641            Self::Continue => "continue",
642            Self::Endloop => "endloop",
643            Self::Call => "call",
644        }
645    }
646}
647#[derive(Debug, Clone, Copy, PartialEq, Eq)]
648#[repr(u8)]
649pub enum SyncOp {
650    Return = 0,
651    Halt = 1,
652    Barrier = 2,
653    FenceAcquire = 3,
654    FenceRelease = 4,
655    FenceAcqRel = 5,
656    Wait = 6,
657    Nop = 7,
658}
659impl SyncOp {
660    #[must_use]
661    pub fn from_u8(value: u8) -> Option<Self> {
662        match value {
663            0 => Some(Self::Return),
664            1 => Some(Self::Halt),
665            2 => Some(Self::Barrier),
666            3 => Some(Self::FenceAcquire),
667            4 => Some(Self::FenceRelease),
668            5 => Some(Self::FenceAcqRel),
669            6 => Some(Self::Wait),
670            7 => Some(Self::Nop),
671            _ => None,
672        }
673    }
674    #[must_use]
675    pub fn mnemonic(self) -> &'static str {
676        match self {
677            Self::Return => "return",
678            Self::Halt => "halt",
679            Self::Barrier => "barrier",
680            Self::FenceAcquire => "fence_acquire",
681            Self::FenceRelease => "fence_release",
682            Self::FenceAcqRel => "fence_acq_rel",
683            Self::Wait => "wait",
684            Self::Nop => "nop",
685        }
686    }
687}
688#[derive(Debug, Clone, Copy, PartialEq, Eq)]
689#[repr(u8)]
690pub enum MiscOp {
691    Mov = 0,
692    MovImm = 1,
693    MovSr = 2,
694}
695impl MiscOp {
696    #[must_use]
697    pub fn from_u8(value: u8) -> Option<Self> {
698        match value {
699            0 => Some(Self::Mov),
700            1 => Some(Self::MovImm),
701            2 => Some(Self::MovSr),
702            _ => None,
703        }
704    }
705    #[must_use]
706    pub fn mnemonic(self) -> &'static str {
707        match self {
708            Self::Mov | Self::MovSr => "mov",
709            Self::MovImm => "mov_imm",
710        }
711    }
712}
713pub const SPECIAL_REGISTER_NAMES: [&str; 16] = [
714    "sr_thread_id_x",
715    "sr_thread_id_y",
716    "sr_thread_id_z",
717    "sr_wave_id",
718    "sr_lane_id",
719    "sr_workgroup_id_x",
720    "sr_workgroup_id_y",
721    "sr_workgroup_id_z",
722    "sr_workgroup_size_x",
723    "sr_workgroup_size_y",
724    "sr_workgroup_size_z",
725    "sr_grid_size_x",
726    "sr_grid_size_y",
727    "sr_grid_size_z",
728    "sr_wave_width",
729    "sr_num_waves",
730];
731#[must_use]
732pub fn special_register_name(index: u8) -> Option<&'static str> {
733    SPECIAL_REGISTER_NAMES.get(index as usize).copied()
734}
wave_decode/opcodes.rs

wave_decode/
opcodes.rs
