llvm-native-core 0.1.16

//! NVPTX Instruction Information — complete PTX instruction
//! descriptor table covering integer, float, compare, logic, data
//! movement, conversion, control flow, atomics, texture, and
//! special operations.
//!
//! Clean-room behavioral reconstruction from the NVIDIA PTX ISA
//! Reference Manual. Zero LLVM source code consultation.

use std::collections::HashMap;

// ============================================================================
// NvptxOpcode — PTX instruction opcodes
// ============================================================================

/// NVPTX instruction opcodes. Each variant corresponds to a PTX
/// instruction mnemonic.  Values are used as `MachineInstr.opcode`.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
#[repr(u32)]
pub enum NvptxOpcode {
    // Integer arithmetic
    Add = 20000,
    Sub = 20001,
    Mul = 20002,
    Mad = 20003,
    Div = 20004,
    Rem = 20005,
    Abs = 20006,
    Neg = 20007,
    Min = 20008,
    Max = 20009,
    Popc = 20010,
    Clz = 20011,
    Bfind = 20012,
    Brev = 20013,
    Bfe = 20014,
    Bfi = 20015,
    Sad = 20016,
    Selp = 20017,

    // Floating-point arithmetic
    FAdd = 20020,
    FSub = 20021,
    FMul = 20022,
    FMad = 20023,
    Fma = 20024,
    FDiv = 20025,
    FAbs = 20026,
    FNeg = 20027,
    FMin = 20028,
    FMax = 20029,
    Rcp = 20030,
    Sqrt = 20031,
    Rsqrt = 20032,
    Sin = 20033,
    Cos = 20034,
    Lg2 = 20035,
    Ex2 = 20036,
    Tanh = 20037,

    // Compare (setp)
    SetpEq = 20040,
    SetpNe = 20041,
    SetpLt = 20042,
    SetpLe = 20043,
    SetpGt = 20044,
    SetpGe = 20045,
    SetpLo = 20046,
    SetpLs = 20047,
    SetpHi = 20048,
    SetpHs = 20049,
    SetpEqu = 20050,
    SetpNeu = 20051,
    SetpLtu = 20052,
    SetpLeu = 20053,
    SetpGtu = 20054,
    SetpGeu = 20055,
    SetpNum = 20056,
    SetpNan = 20057,

    // Logic
    And = 20060,
    Or = 20061,
    Xor = 20062,
    Not = 20063,
    Shl = 20064,
    Shr = 20065,
    Cnot = 20066,

    // Data movement (loads/stores)
    Ld = 20070,
    St = 20071,
    LdGlobal = 20072,
    StGlobal = 20073,
    LdShared = 20074,
    StShared = 20075,
    LdLocal = 20076,
    StLocal = 20077,
    LdParam = 20078,
    StParam = 20079,
    LdConst = 20080,
    LdGlobalNc = 20081,
    LdGlobalCg = 20082,
    LdGlobalCa = 20083,
    Cvta = 20084,
    Mov = 20085,

    // Conversion
    Cvt = 20090,
    CvtRp = 20091,
    CvtRn = 20092,
    CvtRz = 20093,
    CvtRm = 20094,

    // Control flow
    Bra = 20100,
    Call = 20101,
    Ret = 20102,
    Exit = 20103,
    Brkpt = 20104,
    Bar = 20105,
    BarSync = 20106,
    BarArrive = 20107,
    Membar = 20108,

    // Atomics
    AtomAdd = 20110,
    AtomSub = 20111,
    AtomExch = 20112,
    AtomCas = 20113,
    AtomMin = 20114,
    AtomMax = 20115,
    AtomAnd = 20116,
    AtomOr = 20117,
    AtomXor = 20118,

    // Reductions
    RedAdd = 20120,
    RedSub = 20121,
    RedMin = 20122,
    RedMax = 20123,
    RedAnd = 20124,
    RedOr = 20125,
    RedXor = 20126,

    // Vote
    VoteAll = 20130,
    VoteAny = 20131,
    VoteUni = 20132,
    VoteBallot = 20133,

    // Texture / Surface
    Tex = 20140,
    Tld4 = 20141,
    Txq = 20142,
    Suld = 20143,
    Sust = 20144,
    Sured = 20145,
    Suq = 20146,

    // Special
    Isspacep = 20150,
}

impl NvptxOpcode {
    pub fn as_u32(self) -> u32 {
        self as u32
    }

    pub fn is_branch(self) -> bool {
        matches!(self, NvptxOpcode::Bra | NvptxOpcode::Call)
    }

    pub fn is_return(self) -> bool {
        matches!(self, NvptxOpcode::Ret | NvptxOpcode::Exit)
    }

    pub fn is_load(self) -> bool {
        matches!(
            self,
            NvptxOpcode::Ld
                | NvptxOpcode::LdGlobal
                | NvptxOpcode::LdShared
                | NvptxOpcode::LdLocal
                | NvptxOpcode::LdParam
                | NvptxOpcode::LdConst
                | NvptxOpcode::LdGlobalNc
                | NvptxOpcode::LdGlobalCg
                | NvptxOpcode::LdGlobalCa
                | NvptxOpcode::Tex
                | NvptxOpcode::Tld4
                | NvptxOpcode::Suld
        )
    }

    pub fn is_store(self) -> bool {
        matches!(
            self,
            NvptxOpcode::St
                | NvptxOpcode::StGlobal
                | NvptxOpcode::StShared
                | NvptxOpcode::StLocal
                | NvptxOpcode::StParam
                | NvptxOpcode::Sust
        )
    }

    pub fn is_atomic(self) -> bool {
        matches!(
            self,
            NvptxOpcode::AtomAdd
                | NvptxOpcode::AtomSub
                | NvptxOpcode::AtomExch
                | NvptxOpcode::AtomCas
                | NvptxOpcode::AtomMin
                | NvptxOpcode::AtomMax
                | NvptxOpcode::AtomAnd
                | NvptxOpcode::AtomOr
                | NvptxOpcode::AtomXor
        )
    }
}

// ============================================================================
// NvptxOperandType
// ============================================================================

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum NvptxOperandType {
    /// General-purpose virtual register.
    Reg,
    /// Predicate register.
    Pred,
    /// Immediate integer value.
    Imm,
    /// Address space qualifier (e.g., .global, .shared).
    AddrSpace,
    /// Type suffix for conversion/load/store (e.g., .s32, .f32).
    TypeQual,
    /// Label (branch target).
    Label,
    /// Function name for call.
    FuncName,
}

impl NvptxOperandType {
    pub fn is_register(self) -> bool {
        matches!(self, NvptxOperandType::Reg | NvptxOperandType::Pred)
    }

    pub fn is_immediate(self) -> bool {
        matches!(self, NvptxOperandType::Imm)
    }
}

// ============================================================================
// NvptxInstrDesc
// ============================================================================

/// Complete metadata for one PTX instruction.
#[derive(Debug, Clone)]
pub struct NvptxInstrDesc {
    pub opcode: NvptxOpcode,
    pub mnemonic: &'static str,
    pub num_operands: u8,
    pub is_terminator: bool,
    pub is_branch: bool,
    pub is_call: bool,
    pub is_return: bool,
    pub is_compare: bool,
    pub is_load: bool,
    pub is_store: bool,
    pub is_atomic: bool,
    pub has_side_effects: bool,
    pub operand_types: &'static [NvptxOperandType],
}

// ============================================================================
// Static operand type slices
// ============================================================================

const REG_REG_REG: &[NvptxOperandType] = &[
    NvptxOperandType::Reg,
    NvptxOperandType::Reg,
    NvptxOperandType::Reg,
];
const REG_REG_IMM: &[NvptxOperandType] = &[
    NvptxOperandType::Reg,
    NvptxOperandType::Reg,
    NvptxOperandType::Imm,
];
const REG_REG: &[NvptxOperandType] = &[NvptxOperandType::Reg, NvptxOperandType::Reg];
const REG: &[NvptxOperandType] = &[NvptxOperandType::Reg];
const PRED_REG_REG: &[NvptxOperandType] = &[
    NvptxOperandType::Pred,
    NvptxOperandType::Reg,
    NvptxOperandType::Reg,
];
const LABEL: &[NvptxOperandType] = &[NvptxOperandType::Label];
const FUNC: &[NvptxOperandType] = &[NvptxOperandType::FuncName];
const NO_OPS: &[NvptxOperandType] = &[];

// ============================================================================
// NvptxInstrInfo
// ============================================================================

/// Table of PTX instruction descriptors, indexed by `NvptxOpcode`.
pub struct NvptxInstrInfo {
    descriptors: HashMap<NvptxOpcode, NvptxInstrDesc>,
    by_mnemonic: HashMap<String, NvptxOpcode>,
}

impl NvptxInstrInfo {
    pub fn new() -> Self {
        let mut info = NvptxInstrInfo {
            descriptors: HashMap::with_capacity(128),
            by_mnemonic: HashMap::with_capacity(128),
        };
        info.init_descriptors();
        info
    }

    fn add(&mut self, desc: NvptxInstrDesc) {
        self.by_mnemonic
            .insert(desc.mnemonic.to_string(), desc.opcode);
        self.descriptors.insert(desc.opcode, desc);
    }

    fn init_descriptors(&mut self) {
        use NvptxOpcode::*;
        use NvptxOperandType::*;

        // Integer arithmetic
        self.add(instr(Add, "add", 3, false, REG_REG_REG));
        self.add(instr(Sub, "sub", 3, false, REG_REG_REG));
        self.add(instr(Mul, "mul", 3, false, REG_REG_REG));
        self.add(instr(Mad, "mad", 4, false, &[Reg, Reg, Reg, Reg]));
        self.add(instr(Div, "div", 3, false, REG_REG_REG));
        self.add(instr(Rem, "rem", 3, false, REG_REG_REG));
        self.add(instr(Abs, "abs", 2, false, REG_REG));
        self.add(instr(Neg, "neg", 2, false, REG_REG));
        self.add(instr(Min, "min", 3, false, REG_REG_REG));
        self.add(instr(Max, "max", 3, false, REG_REG_REG));
        self.add(instr(Popc, "popc", 2, false, REG_REG));
        self.add(instr(Clz, "clz", 2, false, REG_REG));
        self.add(instr(Bfind, "bfind", 2, false, REG_REG));
        self.add(instr(Brev, "brev", 2, false, REG_REG));
        self.add(instr(Bfe, "bfe", 4, false, &[Reg, Reg, Reg, Reg]));
        self.add(instr(Bfi, "bfi", 4, false, &[Reg, Reg, Reg, Reg]));
        self.add(instr(Sad, "sad", 4, false, &[Reg, Reg, Reg, Reg]));
        self.add(instr_side(Selp, "selp", 3, false, REG_REG_REG));

        // Float
        self.add(instr(FAdd, "add", 3, false, REG_REG_REG));
        self.add(instr(FSub, "sub", 3, false, REG_REG_REG));
        self.add(instr(FMul, "mul", 3, false, REG_REG_REG));
        self.add(instr(FMad, "mad", 4, false, &[Reg, Reg, Reg, Reg]));
        self.add(instr(Fma, "fma", 4, false, &[Reg, Reg, Reg, Reg]));
        self.add(instr(FDiv, "div", 3, false, REG_REG_REG));
        self.add(instr(FAbs, "abs", 2, false, REG_REG));
        self.add(instr(FNeg, "neg", 2, false, REG_REG));
        self.add(instr(FMin, "min", 3, false, REG_REG_REG));
        self.add(instr(FMax, "max", 3, false, REG_REG_REG));
        self.add(instr(Rcp, "rcp", 2, false, REG_REG));
        self.add(instr(Sqrt, "sqrt", 2, false, REG_REG));
        self.add(instr(Rsqrt, "rsqrt", 2, false, REG_REG));
        self.add(instr_side(Sin, "sin", 2, false, REG_REG));
        self.add(instr_side(Cos, "cos", 2, false, REG_REG));
        self.add(instr_side(Lg2, "lg2", 2, false, REG_REG));
        self.add(instr_side(Ex2, "ex2", 2, false, REG_REG));
        self.add(instr_side(Tanh, "tanh", 2, false, REG_REG));

        // Compare
        self.add(instr_comp(SetpEq, "setp.eq", 2, PRED_REG_REG));
        self.add(instr_comp(SetpNe, "setp.ne", 2, PRED_REG_REG));
        self.add(instr_comp(SetpLt, "setp.lt", 2, PRED_REG_REG));
        self.add(instr_comp(SetpLe, "setp.le", 2, PRED_REG_REG));
        self.add(instr_comp(SetpGt, "setp.gt", 2, PRED_REG_REG));
        self.add(instr_comp(SetpGe, "setp.ge", 2, PRED_REG_REG));
        self.add(instr_comp(SetpLo, "setp.lo", 2, PRED_REG_REG));
        self.add(instr_comp(SetpLs, "setp.ls", 2, PRED_REG_REG));
        self.add(instr_comp(SetpHi, "setp.hi", 2, PRED_REG_REG));
        self.add(instr_comp(SetpHs, "setp.hs", 2, PRED_REG_REG));
        self.add(instr_comp(SetpEqu, "setp.equ", 2, PRED_REG_REG));
        self.add(instr_comp(SetpNeu, "setp.neu", 2, PRED_REG_REG));
        self.add(instr_comp(SetpLtu, "setp.ltu", 2, PRED_REG_REG));
        self.add(instr_comp(SetpLeu, "setp.leu", 2, PRED_REG_REG));
        self.add(instr_comp(SetpGtu, "setp.gtu", 2, PRED_REG_REG));
        self.add(instr_comp(SetpGeu, "setp.geu", 2, PRED_REG_REG));
        self.add(instr_comp(SetpNum, "setp.num", 2, PRED_REG_REG));
        self.add(instr_comp(SetpNan, "setp.nan", 2, PRED_REG_REG));

        // Logic
        self.add(instr(And, "and", 3, false, REG_REG_REG));
        self.add(instr(Or, "or", 3, false, REG_REG_REG));
        self.add(instr(Xor, "xor", 3, false, REG_REG_REG));
        self.add(instr(Not, "not", 2, false, REG_REG));
        self.add(instr(Shl, "shl", 3, false, REG_REG_REG));
        self.add(instr(Shr, "shr", 3, false, REG_REG_REG));
        self.add(instr(Cnot, "cnot", 2, false, REG_REG));

        // Data movement — loads
        self.add(instr_load(Ld, "ld"));
        self.add(instr_load(LdGlobal, "ld.global"));
        self.add(instr_load(LdShared, "ld.shared"));
        self.add(instr_load(LdLocal, "ld.local"));
        self.add(instr_load(LdParam, "ld.param"));
        self.add(instr_load(LdConst, "ld.const"));
        self.add(instr_load(LdGlobalNc, "ld.global.nc"));
        self.add(instr_load(LdGlobalCg, "ld.global.cg"));
        self.add(instr_load(LdGlobalCa, "ld.global.ca"));
        // stores
        self.add(instr_store(St, "st"));
        self.add(instr_store(StGlobal, "st.global"));
        self.add(instr_store(StShared, "st.shared"));
        self.add(instr_store(StLocal, "st.local"));
        self.add(instr_store(StParam, "st.param"));

        // Mov and cvta
        self.add(instr(Mov, "mov", 2, false, REG_REG));
        self.add(instr(Cvta, "cvta", 2, false, REG_REG));

        // Conversion
        self.add(instr(Cvt, "cvt", 2, false, REG_REG));
        self.add(instr(CvtRp, "cvt.rp", 2, false, REG_REG));
        self.add(instr(CvtRn, "cvt.rn", 2, false, REG_REG));
        self.add(instr(CvtRz, "cvt.rz", 2, false, REG_REG));
        self.add(instr(CvtRm, "cvt.rm", 2, false, REG_REG));

        // Control flow
        self.add(term(Bra, "bra", 1, true, LABEL));
        self.add(call_instr(Call, "call"));
        self.add(term_ret(Ret, "ret"));
        self.add(term_ret(Exit, "exit"));
        self.add(instr_side(Brkpt, "brkpt", 0, true, NO_OPS));
        self.add(instr_side(Bar, "bar", 1, false, REG));
        self.add(instr_side(BarSync, "bar.sync", 1, false, REG));
        self.add(instr_side(BarArrive, "bar.arrive", 1, false, REG));
        self.add(instr_side(Membar, "membar", 1, false, REG));

        // Atomics
        self.add(instr_atomic(AtomAdd, "atom.add"));
        self.add(instr_atomic(AtomSub, "atom.sub"));
        self.add(instr_atomic(AtomExch, "atom.exch"));
        self.add(instr_atomic(AtomCas, "atom.cas"));
        self.add(instr_atomic(AtomMin, "atom.min"));
        self.add(instr_atomic(AtomMax, "atom.max"));
        self.add(instr_atomic(AtomAnd, "atom.and"));
        self.add(instr_atomic(AtomOr, "atom.or"));
        self.add(instr_atomic(AtomXor, "atom.xor"));

        // Reductions
        self.add(instr_side(RedAdd, "red.add", 2, false, REG_REG));
        self.add(instr_side(RedSub, "red.sub", 2, false, REG_REG));
        self.add(instr_side(RedMin, "red.min", 2, false, REG_REG));
        self.add(instr_side(RedMax, "red.max", 2, false, REG_REG));
        self.add(instr_side(RedAnd, "red.and", 2, false, REG_REG));
        self.add(instr_side(RedOr, "red.or", 2, false, REG_REG));
        self.add(instr_side(RedXor, "red.xor", 2, false, REG_REG));

        // Vote
        self.add(instr_side(VoteAll, "vote.all", 1, false, REG));
        self.add(instr_side(VoteAny, "vote.any", 1, false, REG));
        self.add(instr_side(VoteUni, "vote.uni", 1, false, REG));
        self.add(instr_side(VoteBallot, "vote.ballot", 1, false, REG));

        // Texture / Surface
        self.add(instr_side(Tex, "tex", 2, false, REG_REG));
        self.add(instr_side(Tld4, "tld4", 2, false, REG_REG));
        self.add(instr_side(Txq, "txq", 1, false, REG));
        self.add(instr_side(Suld, "suld", 2, false, REG_REG));
        self.add(instr_side(Sust, "sust", 2, false, REG_REG));
        self.add(instr_side(Sured, "sured", 2, false, REG_REG));
        self.add(instr_side(Suq, "suq", 1, false, REG));

        // Special
        self.add(instr_side(Isspacep, "isspacep", 2, false, REG_REG));
    }

    pub fn get(&self, opcode: NvptxOpcode) -> Option<&NvptxInstrDesc> {
        self.descriptors.get(&opcode)
    }

    pub fn get_mnemonic(&self, opcode: NvptxOpcode) -> &'static str {
        self.descriptors
            .get(&opcode)
            .map(|d| d.mnemonic)
            .unwrap_or("INVALID")
    }

    pub fn find_by_mnemonic(&self, mnemonic: &str) -> Option<NvptxOpcode> {
        self.by_mnemonic.get(mnemonic).copied()
    }

    pub fn is_terminator(&self, opcode: NvptxOpcode) -> bool {
        self.descriptors
            .get(&opcode)
            .map(|d| d.is_terminator)
            .unwrap_or(false)
    }

    pub fn is_branch(&self, opcode: NvptxOpcode) -> bool {
        self.descriptors
            .get(&opcode)
            .map(|d| d.is_branch)
            .unwrap_or(false)
    }

    pub fn may_load(&self, opcode: NvptxOpcode) -> bool {
        self.descriptors
            .get(&opcode)
            .map(|d| d.is_load)
            .unwrap_or(false)
    }

    pub fn may_store(&self, opcode: NvptxOpcode) -> bool {
        self.descriptors
            .get(&opcode)
            .map(|d| d.is_store)
            .unwrap_or(false)
    }

    /// Static mnemonic lookup (doesn't require an NvptxInstrInfo instance).
    pub fn get_mnemonic_static(opcode: NvptxOpcode) -> &'static str {
        match opcode {
            NvptxOpcode::Add => "add",
            NvptxOpcode::Sub => "sub",
            NvptxOpcode::Mul => "mul",
            NvptxOpcode::Mad => "mad",
            NvptxOpcode::Div => "div",
            NvptxOpcode::Rem => "rem",
            NvptxOpcode::Abs => "abs",
            NvptxOpcode::Neg => "neg",
            NvptxOpcode::Min => "min",
            NvptxOpcode::Max => "max",
            NvptxOpcode::Popc => "popc",
            NvptxOpcode::Clz => "clz",
            NvptxOpcode::Bfind => "bfind",
            NvptxOpcode::Brev => "brev",
            NvptxOpcode::Bfe => "bfe",
            NvptxOpcode::Bfi => "bfi",
            NvptxOpcode::Sad => "sad",
            NvptxOpcode::Selp => "selp",
            NvptxOpcode::FAdd => "add",
            NvptxOpcode::FSub => "sub",
            NvptxOpcode::FMul => "mul",
            NvptxOpcode::FMad => "mad",
            NvptxOpcode::Fma => "fma",
            NvptxOpcode::FDiv => "div",
            NvptxOpcode::FAbs => "abs",
            NvptxOpcode::FNeg => "neg",
            NvptxOpcode::FMin => "min",
            NvptxOpcode::FMax => "max",
            NvptxOpcode::Rcp => "rcp",
            NvptxOpcode::Sqrt => "sqrt",
            NvptxOpcode::Rsqrt => "rsqrt",
            NvptxOpcode::Sin => "sin",
            NvptxOpcode::Cos => "cos",
            NvptxOpcode::Lg2 => "lg2",
            NvptxOpcode::Ex2 => "ex2",
            NvptxOpcode::Tanh => "tanh",
            NvptxOpcode::SetpEq => "setp.eq",
            NvptxOpcode::SetpNe => "setp.ne",
            NvptxOpcode::SetpLt => "setp.lt",
            NvptxOpcode::SetpLe => "setp.le",
            NvptxOpcode::SetpGt => "setp.gt",
            NvptxOpcode::SetpGe => "setp.ge",
            NvptxOpcode::SetpLo => "setp.lo",
            NvptxOpcode::SetpLs => "setp.ls",
            NvptxOpcode::SetpHi => "setp.hi",
            NvptxOpcode::SetpHs => "setp.hs",
            NvptxOpcode::SetpEqu => "setp.equ",
            NvptxOpcode::SetpNeu => "setp.neu",
            NvptxOpcode::SetpLtu => "setp.ltu",
            NvptxOpcode::SetpLeu => "setp.leu",
            NvptxOpcode::SetpGtu => "setp.gtu",
            NvptxOpcode::SetpGeu => "setp.geu",
            NvptxOpcode::SetpNum => "setp.num",
            NvptxOpcode::SetpNan => "setp.nan",
            NvptxOpcode::And => "and",
            NvptxOpcode::Or => "or",
            NvptxOpcode::Xor => "xor",
            NvptxOpcode::Not => "not",
            NvptxOpcode::Shl => "shl",
            NvptxOpcode::Shr => "shr",
            NvptxOpcode::Cnot => "cnot",
            NvptxOpcode::Ld => "ld",
            NvptxOpcode::St => "st",
            NvptxOpcode::LdGlobal => "ld.global",
            NvptxOpcode::StGlobal => "st.global",
            NvptxOpcode::LdShared => "ld.shared",
            NvptxOpcode::StShared => "st.shared",
            NvptxOpcode::LdLocal => "ld.local",
            NvptxOpcode::StLocal => "st.local",
            NvptxOpcode::LdParam => "ld.param",
            NvptxOpcode::StParam => "st.param",
            NvptxOpcode::LdConst => "ld.const",
            NvptxOpcode::LdGlobalNc => "ld.global.nc",
            NvptxOpcode::LdGlobalCg => "ld.global.cg",
            NvptxOpcode::LdGlobalCa => "ld.global.ca",
            NvptxOpcode::Cvta => "cvta",
            NvptxOpcode::Mov => "mov",
            NvptxOpcode::Cvt => "cvt",
            NvptxOpcode::CvtRp => "cvt.rp",
            NvptxOpcode::CvtRn => "cvt.rn",
            NvptxOpcode::CvtRz => "cvt.rz",
            NvptxOpcode::CvtRm => "cvt.rm",
            NvptxOpcode::Bra => "bra",
            NvptxOpcode::Call => "call",
            NvptxOpcode::Ret => "ret",
            NvptxOpcode::Exit => "exit",
            NvptxOpcode::Brkpt => "brkpt",
            NvptxOpcode::Bar => "bar",
            NvptxOpcode::BarSync => "bar.sync",
            NvptxOpcode::BarArrive => "bar.arrive",
            NvptxOpcode::Membar => "membar",
            NvptxOpcode::AtomAdd => "atom.add",
            NvptxOpcode::AtomSub => "atom.sub",
            NvptxOpcode::AtomExch => "atom.exch",
            NvptxOpcode::AtomCas => "atom.cas",
            NvptxOpcode::AtomMin => "atom.min",
            NvptxOpcode::AtomMax => "atom.max",
            NvptxOpcode::AtomAnd => "atom.and",
            NvptxOpcode::AtomOr => "atom.or",
            NvptxOpcode::AtomXor => "atom.xor",
            NvptxOpcode::RedAdd => "red.add",
            NvptxOpcode::RedSub => "red.sub",
            NvptxOpcode::RedMin => "red.min",
            NvptxOpcode::RedMax => "red.max",
            NvptxOpcode::RedAnd => "red.and",
            NvptxOpcode::RedOr => "red.or",
            NvptxOpcode::RedXor => "red.xor",
            NvptxOpcode::VoteAll => "vote.all",
            NvptxOpcode::VoteAny => "vote.any",
            NvptxOpcode::VoteUni => "vote.uni",
            NvptxOpcode::VoteBallot => "vote.ballot",
            NvptxOpcode::Tex => "tex",
            NvptxOpcode::Tld4 => "tld4",
            NvptxOpcode::Txq => "txq",
            NvptxOpcode::Suld => "suld",
            NvptxOpcode::Sust => "sust",
            NvptxOpcode::Sured => "sured",
            NvptxOpcode::Suq => "suq",
            NvptxOpcode::Isspacep => "isspacep",
        }
    }
}

impl Default for NvptxInstrInfo {
    fn default() -> Self {
        Self::new()
    }
}

// ============================================================================
// Construction helpers
// ============================================================================

fn instr(
    opcode: NvptxOpcode,
    mnemonic: &'static str,
    num_operands: u8,
    has_side_effects: bool,
    operand_types: &'static [NvptxOperandType],
) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands,
        is_terminator: false,
        is_branch: false,
        is_call: false,
        is_return: false,
        is_compare: false,
        is_load: false,
        is_store: false,
        is_atomic: false,
        has_side_effects,
        operand_types,
    }
}

fn instr_side(
    opcode: NvptxOpcode,
    mnemonic: &'static str,
    num_operands: u8,
    has_side_effects: bool,
    operand_types: &'static [NvptxOperandType],
) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands,
        is_terminator: false,
        is_branch: false,
        is_call: false,
        is_return: false,
        is_compare: false,
        is_load: false,
        is_store: false,
        is_atomic: false,
        has_side_effects,
        operand_types,
    }
}

fn instr_comp(
    opcode: NvptxOpcode,
    mnemonic: &'static str,
    num_operands: u8,
    operand_types: &'static [NvptxOperandType],
) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands,
        is_terminator: false,
        is_branch: false,
        is_call: false,
        is_return: false,
        is_compare: true,
        is_load: false,
        is_store: false,
        is_atomic: false,
        has_side_effects: false,
        operand_types,
    }
}

fn instr_load(opcode: NvptxOpcode, mnemonic: &'static str) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands: 2,
        is_terminator: false,
        is_branch: false,
        is_call: false,
        is_return: false,
        is_compare: false,
        is_load: true,
        is_store: false,
        is_atomic: false,
        has_side_effects: true,
        operand_types: REG_REG,
    }
}

fn instr_store(opcode: NvptxOpcode, mnemonic: &'static str) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands: 2,
        is_terminator: false,
        is_branch: false,
        is_call: false,
        is_return: false,
        is_compare: false,
        is_load: false,
        is_store: true,
        is_atomic: false,
        has_side_effects: true,
        operand_types: REG_REG,
    }
}

fn instr_atomic(opcode: NvptxOpcode, mnemonic: &'static str) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands: 3,
        is_terminator: false,
        is_branch: false,
        is_call: false,
        is_return: false,
        is_compare: false,
        is_load: false,
        is_store: false,
        is_atomic: true,
        has_side_effects: true,
        operand_types: REG_REG_REG,
    }
}

fn term(
    opcode: NvptxOpcode,
    mnemonic: &'static str,
    num_operands: u8,
    has_side_effects: bool,
    operand_types: &'static [NvptxOperandType],
) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands,
        is_terminator: true,
        is_branch: true,
        is_call: false,
        is_return: false,
        is_compare: false,
        is_load: false,
        is_store: false,
        is_atomic: false,
        has_side_effects,
        operand_types,
    }
}

fn term_ret(opcode: NvptxOpcode, mnemonic: &'static str) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands: 0,
        is_terminator: true,
        is_branch: false,
        is_call: false,
        is_return: true,
        is_compare: false,
        is_load: false,
        is_store: false,
        is_atomic: false,
        has_side_effects: true,
        operand_types: NO_OPS,
    }
}

fn call_instr(opcode: NvptxOpcode, mnemonic: &'static str) -> NvptxInstrDesc {
    NvptxInstrDesc {
        opcode,
        mnemonic,
        num_operands: 1,
        is_terminator: false,
        is_branch: false,
        is_call: true,
        is_return: false,
        is_compare: false,
        is_load: false,
        is_store: false,
        is_atomic: false,
        has_side_effects: true,
        operand_types: FUNC,
    }
}

#[cfg(test)]
mod tests {
    use super::*;

    #[test]
    fn test_instr_info_get() {
        let info = NvptxInstrInfo::new();
        let desc = info.get(NvptxOpcode::Add).unwrap();
        assert_eq!(desc.mnemonic, "add");
        assert_eq!(desc.num_operands, 3);
    }

    #[test]
    fn test_find_by_mnemonic() {
        let info = NvptxInstrInfo::new();
        assert_eq!(info.find_by_mnemonic("popc"), Some(NvptxOpcode::Popc));
        assert_eq!(info.find_by_mnemonic("brev"), Some(NvptxOpcode::Brev));
        assert_eq!(info.find_by_mnemonic("bra"), Some(NvptxOpcode::Bra));
        assert_eq!(info.find_by_mnemonic("ret"), Some(NvptxOpcode::Ret));
    }

    #[test]
    fn test_terminators() {
        let info = NvptxInstrInfo::new();
        assert!(info.is_terminator(NvptxOpcode::Bra));
        assert!(info.is_terminator(NvptxOpcode::Ret));
        assert!(!info.is_terminator(NvptxOpcode::Add));
    }

    #[test]
    fn test_load_store() {
        let info = NvptxInstrInfo::new();
        assert!(info.may_load(NvptxOpcode::Ld));
        assert!(info.may_load(NvptxOpcode::LdGlobal));
        assert!(!info.may_load(NvptxOpcode::St));
        assert!(info.may_store(NvptxOpcode::St));
        assert!(info.may_store(NvptxOpcode::StGlobal));
        assert!(!info.may_store(NvptxOpcode::Ld));
    }

    #[test]
    fn test_compare_instructions() {
        let info = NvptxInstrInfo::new();
        let desc = info.get(NvptxOpcode::SetpEq).unwrap();
        assert!(desc.is_compare);
    }

    #[test]
    fn test_atomic_instructions() {
        let info = NvptxInstrInfo::new();
        let desc = info.get(NvptxOpcode::AtomAdd).unwrap();
        assert!(desc.is_atomic);
    }
}