llvm-native-core 0.1.15

//! NVPTX Complete Instruction Selection — compute capability specific patterns,
//! full PTX instruction set, and address space lowering.
//!
//! Clean-room behavioral reconstruction from the NVIDIA PTX ISA Reference Manual.
//! Zero LLVM source code consultation.
//!
//! # PTX Instruction Categories
//! - Integer arithmetic: add, sub, mul, mad, div, rem, abs, neg, min, max
//! - Float arithmetic: add.f32/f64, sub, mul, fma, div, rcp, sqrt, rsqrt
//! - Comparison: setp.eq/ne/lt/le/gt/ge/lo/ls/hi/hs/num/nan
//! - Bitwise: and, or, xor, not, shl, shr, cnot, popc, clz, bfind, brev, bfe, bfi
//! - Memory: ld (global/local/shared/const), st, atom
//! - Synchronization: bar.sync, membar, atom.global
//! - Texture/Surface: tex, tld4, suq, suld, sust, surd
//! - Special: vote, shfl, prmt, ldmatrix, mma, stmatrix
//! - Address spaces: .global, .local, .shared, .const, .param

use super::nvptx_instr_info::NvptxOpcode;
use crate::codegen::*;
use crate::opcode::Opcode;
use crate::value::Value;
use std::collections::HashMap;

// ============================================================================
// Compute Capability Types
// ============================================================================

/// NVIDIA compute capability version
#[derive(Debug, Clone, Copy, PartialEq, Eq, PartialOrd, Ord, Hash)]
pub enum ComputeCapability {
    Sm20, // Fermi
    Sm21,
    Sm30, // Kepler
    Sm32,
    Sm35, // Kepler refresh
    Sm37,
    Sm50, // Maxwell
    Sm52,
    Sm53,
    Sm60, // Pascal
    Sm61,
    Sm62,
    Sm70, // Volta
    Sm72,
    Sm75, // Turing
    Sm80, // Ampere
    Sm86,
    Sm87,
    Sm89,
    Sm90,  // Hopper
    Sm90a, // Hopper with sm_90a features
}

impl ComputeCapability {
    pub fn from_pair(major: u32, minor: u32) -> Option<Self> {
        match (major, minor) {
            (2, 0) => Some(ComputeCapability::Sm20),
            (2, 1) => Some(ComputeCapability::Sm21),
            (3, 0) => Some(ComputeCapability::Sm30),
            (3, 2) => Some(ComputeCapability::Sm32),
            (3, 5) => Some(ComputeCapability::Sm35),
            (3, 7) => Some(ComputeCapability::Sm37),
            (5, 0) => Some(ComputeCapability::Sm50),
            (5, 2) => Some(ComputeCapability::Sm52),
            (5, 3) => Some(ComputeCapability::Sm53),
            (6, 0) => Some(ComputeCapability::Sm60),
            (6, 1) => Some(ComputeCapability::Sm61),
            (6, 2) => Some(ComputeCapability::Sm62),
            (7, 0) => Some(ComputeCapability::Sm70),
            (7, 2) => Some(ComputeCapability::Sm72),
            (7, 5) => Some(ComputeCapability::Sm75),
            (8, 0) => Some(ComputeCapability::Sm80),
            (8, 6) => Some(ComputeCapability::Sm86),
            (8, 7) => Some(ComputeCapability::Sm87),
            (8, 9) => Some(ComputeCapability::Sm89),
            (9, 0) => Some(ComputeCapability::Sm90),
            _ => None,
        }
    }

    pub fn from_name(name: &str) -> Option<Self> {
        let lower = name.to_lowercase();
        if lower.starts_with("sm_") {
            let parts: Vec<&str> = lower[3..].split('a').collect();
            let nums: Vec<&str> = parts[0].split('_').collect();
            if nums.len() == 2 {
                let major: u32 = nums[0].parse().ok()?;
                let minor: u32 = nums[1].parse().ok()?;
                let cap = Self::from_pair(major, minor)?;
                if parts.len() > 1 && cap == ComputeCapability::Sm90 {
                    return Some(ComputeCapability::Sm90a);
                }
                return Some(cap);
            }
        }
        None
    }

    pub fn major(self) -> u32 {
        match self {
            ComputeCapability::Sm20 | ComputeCapability::Sm21 => 2,
            ComputeCapability::Sm30
            | ComputeCapability::Sm32
            | ComputeCapability::Sm35
            | ComputeCapability::Sm37 => 3,
            ComputeCapability::Sm50 | ComputeCapability::Sm52 | ComputeCapability::Sm53 => 5,
            ComputeCapability::Sm60 | ComputeCapability::Sm61 | ComputeCapability::Sm62 => 6,
            ComputeCapability::Sm70 | ComputeCapability::Sm72 | ComputeCapability::Sm75 => 7,
            ComputeCapability::Sm80
            | ComputeCapability::Sm86
            | ComputeCapability::Sm87
            | ComputeCapability::Sm89 => 8,
            ComputeCapability::Sm90 | ComputeCapability::Sm90a => 9,
        }
    }

    pub fn minor(self) -> u32 {
        match self {
            ComputeCapability::Sm20 => 0,
            ComputeCapability::Sm21 => 1,
            ComputeCapability::Sm30 => 0,
            ComputeCapability::Sm32 => 2,
            ComputeCapability::Sm35 => 5,
            ComputeCapability::Sm37 => 7,
            ComputeCapability::Sm50 => 0,
            ComputeCapability::Sm52 => 2,
            ComputeCapability::Sm53 => 3,
            ComputeCapability::Sm60 => 0,
            ComputeCapability::Sm61 => 1,
            ComputeCapability::Sm62 => 2,
            ComputeCapability::Sm70 => 0,
            ComputeCapability::Sm72 => 2,
            ComputeCapability::Sm75 => 5,
            ComputeCapability::Sm80 => 0,
            ComputeCapability::Sm86 => 6,
            ComputeCapability::Sm87 => 7,
            ComputeCapability::Sm89 => 9,
            ComputeCapability::Sm90 | ComputeCapability::Sm90a => 0,
        }
    }

    pub fn supports_mma(self) -> bool {
        self >= ComputeCapability::Sm70
    }

    pub fn supports_dp_fma(self) -> bool {
        self >= ComputeCapability::Sm60
    }

    pub fn supports_tensor_core(self) -> bool {
        self >= ComputeCapability::Sm70
    }

    pub fn supports_sparse_tensor_core(self) -> bool {
        self >= ComputeCapability::Sm80
    }

    pub fn supports_ldmatrix(self) -> bool {
        self >= ComputeCapability::Sm75
    }

    pub fn warp_size(self) -> u32 {
        32
    }
    pub fn num_warps_per_sm(self) -> u32 {
        match self.major() {
            2 => 48,
            3..=5 => 64,
            6 => 64,
            7 => 64,
            8 => 64,
            9 => 64,
            _ => 32,
        }
    }

    pub fn num_registers_per_sm(self) -> u32 {
        match self.major() {
            2 => 32768,
            3 => 65536,
            5 => 65536,
            6 => 65536,
            7 => 65536,
            8 => 65536,
            9 => 65536,
            _ => 32768,
        }
    }
}

// ============================================================================
// PTX Address Space
// ============================================================================

/// PTX address spaces for memory operations
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum PtxAddrSpace {
    Global,
    Local,
    Shared,
    Const,
    Param,
    Generic,
    Texture,
    Surface,
}

impl PtxAddrSpace {
    pub fn from_u32(val: u32) -> Option<Self> {
        match val {
            0 => Some(PtxAddrSpace::Generic),
            1 => Some(PtxAddrSpace::Global),
            2 => Some(PtxAddrSpace::Local),
            3 => Some(PtxAddrSpace::Shared),
            4 => Some(PtxAddrSpace::Const),
            5 => Some(PtxAddrSpace::Param),
            _ => None,
        }
    }

    pub fn to_ptx_qualifier(self) -> &'static str {
        match self {
            PtxAddrSpace::Global => ".global",
            PtxAddrSpace::Local => ".local",
            PtxAddrSpace::Shared => ".shared",
            PtxAddrSpace::Const => ".const",
            PtxAddrSpace::Param => ".param",
            PtxAddrSpace::Generic => "",
            PtxAddrSpace::Texture => ".tex",
            PtxAddrSpace::Surface => ".surf",
        }
    }
}

// ============================================================================
// PTX Data Type
// ============================================================================

/// PTX type specifier for instructions
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum PtxType {
    B8,
    B16,
    B32,
    B64,
    U8,
    U16,
    U32,
    U64,
    S8,
    S16,
    S32,
    S64,
    F16,
    F16x2,
    F32,
    F64,
    Pred,
    B128,
}

impl PtxType {
    pub fn to_str(self) -> &'static str {
        match self {
            PtxType::B8 => "b8",
            PtxType::B16 => "b16",
            PtxType::B32 => "b32",
            PtxType::B64 => "b64",
            PtxType::U8 => "u8",
            PtxType::U16 => "u16",
            PtxType::U32 => "u32",
            PtxType::U64 => "u64",
            PtxType::S8 => "s8",
            PtxType::S16 => "s16",
            PtxType::S32 => "s32",
            PtxType::S64 => "s64",
            PtxType::F16 => "f16",
            PtxType::F16x2 => "f16x2",
            PtxType::F32 => "f32",
            PtxType::F64 => "f64",
            PtxType::Pred => "pred",
            PtxType::B128 => "b128",
        }
    }

    pub fn size_bytes(self) -> u32 {
        match self {
            PtxType::B8 | PtxType::U8 | PtxType::S8 => 1,
            PtxType::B16 | PtxType::U16 | PtxType::S16 | PtxType::F16 | PtxType::F16x2 => 2,
            PtxType::B32 | PtxType::U32 | PtxType::S32 | PtxType::F32 => 4,
            PtxType::B64 | PtxType::U64 | PtxType::S64 | PtxType::F64 => 8,
            PtxType::Pred => 1,
            PtxType::B128 => 16,
        }
    }
}

// ============================================================================
// PTX Atomic Operation
// ============================================================================

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum PtxAtomicOp {
    Add,
    Sub,
    Exch,
    Cas,
    Min,
    Max,
    And,
    Or,
    Xor,
    Inc,
    Dec,
}

impl PtxAtomicOp {
    pub fn to_str(self) -> &'static str {
        match self {
            PtxAtomicOp::Add => "add",
            PtxAtomicOp::Sub => "sub",
            PtxAtomicOp::Exch => "exch",
            PtxAtomicOp::Cas => "cas",
            PtxAtomicOp::Min => "min",
            PtxAtomicOp::Max => "max",
            PtxAtomicOp::And => "and",
            PtxAtomicOp::Or => "or",
            PtxAtomicOp::Xor => "xor",
            PtxAtomicOp::Inc => "inc",
            PtxAtomicOp::Dec => "dec",
        }
    }
}

// ============================================================================
// Complete NVPTX Instruction Selector
// ============================================================================

pub struct NvptxFullInstructionSelector {
    pub cc: ComputeCapability,
    pub use_fast_math: bool,
    pub addr_size: u32,
    pub vreg_map: HashMap<usize, VirtReg>,
    pub pred_map: HashMap<usize, VirtReg>,
    pub mbb: MachineBasicBlock,
    pub func_name: String,
    pub default_addr_space: PtxAddrSpace,
}

impl NvptxFullInstructionSelector {
    pub fn new(cc: ComputeCapability) -> Self {
        NvptxFullInstructionSelector {
            cc,
            use_fast_math: false,
            addr_size: 64,
            vreg_map: HashMap::new(),
            pred_map: HashMap::new(),
            mbb: MachineBasicBlock {
                name: String::new(),
                instructions: Vec::new(),
                successors: Vec::new(),
            },
            func_name: String::new(),
            default_addr_space: PtxAddrSpace::Global,
        }
    }

    pub fn with_fast_math(mut self, fm: bool) -> Self {
        self.use_fast_math = fm;
        self
    }

    // ========================================================================
    // Integer arithmetic lowering
    // ========================================================================

    pub fn lower_add(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Add as u32)
    }

    pub fn lower_sub(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Sub as u32)
    }

    pub fn lower_mul(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Mul as u32)
    }

    pub fn lower_mad(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        // mad = a * b + c (fused multiply-add for integers)
        self.lower_four_reg(inst, NvptxOpcode::Mad as u32)
    }

    pub fn lower_sdiv(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Div as u32)
    }

    pub fn lower_udiv(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Div as u32)
    }

    pub fn lower_rem(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Rem as u32)
    }

    pub fn lower_abs(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Abs as u32)
    }

    pub fn lower_neg(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Neg as u32)
    }

    pub fn lower_min(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Min as u32)
    }

    pub fn lower_max(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Max as u32)
    }

    // ========================================================================
    // Float arithmetic lowering
    // ========================================================================

    pub fn lower_fadd(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        if self.use_fast_math {
            self.lower_three_reg(inst, NvptxOpcode::FAdd as u32)
        } else {
            self.lower_three_reg(inst, NvptxOpcode::FAdd as u32)
        }
    }

    pub fn lower_fsub(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::FSub as u32)
    }

    pub fn lower_fmul(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::FMul as u32)
    }

    pub fn lower_fma(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        // FMA: d = a * b + c
        let op = if self.cc.supports_dp_fma() || ty != PtxType::F64 {
            NvptxOpcode::Fma as u32
        } else {
            NvptxOpcode::FMad as u32
        };
        self.lower_four_reg(inst, op)
    }

    pub fn lower_fdiv(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::FDiv as u32)
    }

    pub fn lower_rcp(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Rcp as u32)
    }

    pub fn lower_sqrt(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Sqrt as u32)
    }

    pub fn lower_rsqrt(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Rsqrt as u32)
    }

    pub fn lower_fabs(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::FAbs as u32)
    }

    pub fn lower_fneg(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::FNeg as u32)
    }

    pub fn lower_fmin(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::FMin as u32)
    }

    pub fn lower_fmax(&self, inst: &Value, ty: PtxType) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::FMax as u32)
    }

    // Special math functions
    pub fn lower_sin(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Sin as u32)
    }

    pub fn lower_cos(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Cos as u32)
    }

    pub fn lower_lg2(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Lg2 as u32)
    }

    pub fn lower_ex2(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Ex2 as u32)
    }

    pub fn lower_tanh(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Tanh as u32)
    }

    // ========================================================================
    // Comparison lowering
    // ========================================================================

    pub fn lower_setp_eq(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpEq as u32)
    }

    pub fn lower_setp_ne(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpNe as u32)
    }

    pub fn lower_setp_lt(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpLt as u32)
    }

    pub fn lower_setp_le(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpLe as u32)
    }

    pub fn lower_setp_gt(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpGt as u32)
    }

    pub fn lower_setp_ge(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpGe as u32)
    }

    // Unsigned comparisons
    pub fn lower_setp_lo(&self, inst: &Value) -> MachineInstr {
        // lo = "lower" unsigned
        self.lower_cmp_two(inst, NvptxOpcode::SetpLo as u32)
    }

    pub fn lower_setp_ls(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpLs as u32)
    }

    pub fn lower_setp_hi(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpHi as u32)
    }

    pub fn lower_setp_hs(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpHs as u32)
    }

    pub fn lower_setp_num(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpNum as u32)
    }

    pub fn lower_setp_nan(&self, inst: &Value) -> MachineInstr {
        self.lower_cmp_two(inst, NvptxOpcode::SetpNan as u32)
    }

    // ========================================================================
    // Bitwise lowering
    // ========================================================================

    pub fn lower_and(&self, inst: &Value) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::And as u32)
    }

    pub fn lower_or(&self, inst: &Value) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Or as u32)
    }

    pub fn lower_xor(&self, inst: &Value) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Xor as u32)
    }

    pub fn lower_not(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Not as u32)
    }

    pub fn lower_shl(&self, inst: &Value) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Shl as u32)
    }

    pub fn lower_shr(&self, inst: &Value) -> MachineInstr {
        self.lower_three_reg(inst, NvptxOpcode::Shr as u32)
    }

    pub fn lower_cnot(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Cnot as u32)
    }

    pub fn lower_popc(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Popc as u32)
    }

    pub fn lower_clz(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Clz as u32)
    }

    pub fn lower_bfind(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Bfind as u32)
    }

    pub fn lower_brev(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Brev as u32)
    }

    // ========================================================================
    // Memory lowering (load/store)
    // ========================================================================

    pub fn lower_load(&self, inst: &Value, space: PtxAddrSpace) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let addr: u32 = self.get_vreg_for_operand(inst, 0);
        let op = match space {
            PtxAddrSpace::Global => NvptxOpcode::LdGlobal as u32,
            PtxAddrSpace::Local => NvptxOpcode::LdLocal as u32,
            PtxAddrSpace::Shared => NvptxOpcode::LdShared as u32,
            PtxAddrSpace::Const => NvptxOpcode::LdConst as u32,
            PtxAddrSpace::Param => NvptxOpcode::LdParam as u32,
            PtxAddrSpace::Generic => NvptxOpcode::Ld as u32,
            _ => NvptxOpcode::Ld as u32,
        };
        let mut mi = MachineInstr::new(op);
        mi.push_reg(rd);
        mi.push_reg(addr);
        mi.push_imm(0);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_store(&self, inst: &Value, space: PtxAddrSpace) -> MachineInstr {
        let val: u32 = self.get_vreg_for_operand(inst, 0);
        let addr: u32 = self.get_vreg_for_operand(inst, 1);
        let op = match space {
            PtxAddrSpace::Global => NvptxOpcode::StGlobal as u32,
            PtxAddrSpace::Local => NvptxOpcode::StLocal as u32,
            PtxAddrSpace::Shared => NvptxOpcode::StShared as u32,
            PtxAddrSpace::Param => NvptxOpcode::StParam as u32,
            PtxAddrSpace::Generic => NvptxOpcode::St as u32,
            _ => NvptxOpcode::St as u32,
        };
        let mut mi = MachineInstr::new(op);
        mi.push_reg(val);
        mi.push_reg(addr);
        mi.push_imm(0);
        mi
    }

    pub fn lower_load_global_nc(&self, inst: &Value) -> MachineInstr {
        // Non-coherent global load (faster, skips L1)
        self.lower_two_reg_imm(inst, NvptxOpcode::LdGlobalNc as u32)
    }

    pub fn lower_load_global_cg(&self, inst: &Value) -> MachineInstr {
        // Cache-global (L1 cached)
        self.lower_two_reg_imm(inst, NvptxOpcode::LdGlobalCg as u32)
    }

    pub fn lower_load_global_ca(&self, inst: &Value) -> MachineInstr {
        // Cache-all (L1 and L2 cached)
        self.lower_two_reg_imm(inst, NvptxOpcode::LdGlobalCa as u32)
    }

    // ========================================================================
    // Atomic operations
    // ========================================================================

    pub fn lower_atomic(
        &self,
        inst: &Value,
        op: PtxAtomicOp,
        space: PtxAddrSpace,
        ty: PtxType,
    ) -> MachineInstr {
        self.lower_atomic_inner(inst, op, space)
    }

    fn lower_atomic_inner(
        &self,
        inst: &Value,
        op: PtxAtomicOp,
        _space: PtxAddrSpace,
    ) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let addr: u32 = self.get_vreg_for_operand(inst, 0);
        let val: u32 = self.get_vreg_for_operand(inst, 1);
        // Map to appropriate opcode
        let opcode = match op {
            PtxAtomicOp::Add => NvptxOpcode::AtomAdd as u32,
            PtxAtomicOp::Sub => NvptxOpcode::AtomSub as u32,
            PtxAtomicOp::Exch => NvptxOpcode::AtomExch as u32,
            PtxAtomicOp::Cas => NvptxOpcode::AtomCas as u32,
            PtxAtomicOp::Min => NvptxOpcode::AtomMin as u32,
            PtxAtomicOp::Max => NvptxOpcode::AtomMax as u32,
            PtxAtomicOp::And => NvptxOpcode::AtomAnd as u32,
            PtxAtomicOp::Or => NvptxOpcode::AtomOr as u32,
            PtxAtomicOp::Xor => NvptxOpcode::AtomXor as u32,
            _ => NvptxOpcode::AtomAdd as u32,
        };
        let mut mi = MachineInstr::new(opcode);
        mi.push_reg(rd);
        mi.push_reg(addr);
        mi.push_reg(val);
        if inst.operands.len() >= 3 {
            let cmp: u32 = self.get_vreg_for_operand(inst, 2);
            mi.push_reg(cmp);
        }
        mi.def = Some(rd);
        mi
    }

    // ========================================================================
    // Reduction operations
    // ========================================================================

    pub fn lower_red(&self, inst: &Value, op: PtxAtomicOp) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let val: u32 = self.get_vreg_for_operand(inst, 0);
        let opcode = match op {
            PtxAtomicOp::Add => NvptxOpcode::RedAdd as u32,
            PtxAtomicOp::Sub => NvptxOpcode::RedSub as u32,
            PtxAtomicOp::Min => NvptxOpcode::RedMin as u32,
            PtxAtomicOp::Max => NvptxOpcode::RedMax as u32,
            PtxAtomicOp::And => NvptxOpcode::RedAnd as u32,
            PtxAtomicOp::Or => NvptxOpcode::RedOr as u32,
            PtxAtomicOp::Xor => NvptxOpcode::RedXor as u32,
            _ => NvptxOpcode::RedAdd as u32,
        };
        let mut mi = MachineInstr::new(opcode);
        mi.push_reg(rd);
        mi.push_reg(val);
        mi.def = Some(rd);
        mi
    }

    // ========================================================================
    // Synchronization and control
    // ========================================================================

    pub fn lower_bar_sync(&self, inst: &Value) -> MachineInstr {
        let mut mi = MachineInstr::new(NvptxOpcode::BarSync as u32);
        mi.push_imm(0);
        mi
    }

    pub fn lower_bar_arrive(&self, inst: &Value) -> MachineInstr {
        let mut mi = MachineInstr::new(NvptxOpcode::BarArrive as u32);
        mi.push_imm(0);
        mi
    }

    pub fn lower_membar(&self, level: &str) -> MachineInstr {
        let mut mi = MachineInstr::new(NvptxOpcode::Membar as u32);
        mi.push_imm(match level {
            "cta" => 0,
            "gl" => 1,
            "sys" => 2,
            _ => 0,
        });
        mi
    }

    pub fn lower_bra(&self, target: &str) -> MachineInstr {
        let mut mi = MachineInstr::new(NvptxOpcode::Bra as u32);
        mi.push_label(target);
        mi
    }

    pub fn lower_call(&self, target: &str) -> MachineInstr {
        let mut mi = MachineInstr::new(NvptxOpcode::Call as u32);
        mi.push_label(target);
        mi
    }

    pub fn lower_ret(&self) -> MachineInstr {
        MachineInstr::new(NvptxOpcode::Ret as u32)
    }

    pub fn lower_exit(&self) -> MachineInstr {
        MachineInstr::new(NvptxOpcode::Exit as u32)
    }

    pub fn lower_brkpt(&self) -> MachineInstr {
        MachineInstr::new(NvptxOpcode::Brkpt as u32)
    }

    // ========================================================================
    // Voting operations
    // ========================================================================

    pub fn lower_vote_all(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg_pred(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::VoteAll as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_vote_any(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg_pred(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::VoteAny as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_vote_uni(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg_pred(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::VoteUni as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_vote_ballot(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let mut mi = MachineInstr::new(NvptxOpcode::VoteBallot as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    // ========================================================================
    // Texture/Surface operations
    // ========================================================================

    pub fn lower_tex(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let mut mi = MachineInstr::new(NvptxOpcode::Tex as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_tld4(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let mut mi = MachineInstr::new(NvptxOpcode::Tld4 as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_txq(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let mut mi = MachineInstr::new(NvptxOpcode::Txq as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_suld(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let mut mi = MachineInstr::new(NvptxOpcode::Suld as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_sust(&self, inst: &Value) -> MachineInstr {
        let val: u32 = self.get_vreg_for_operand(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::Sust as u32);
        mi.push_reg(val);
        mi
    }

    pub fn lower_sured(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let mut mi = MachineInstr::new(NvptxOpcode::Sured as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_suq(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let mut mi = MachineInstr::new(NvptxOpcode::Suq as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    // ========================================================================
    // Address space predicate
    // ========================================================================

    pub fn lower_isspacep(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg_pred(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::Isspacep as u32);
        mi.push_reg(rd);
        mi.def = Some(rd);
        mi
    }

    // ========================================================================
    // Conversion operations
    // ========================================================================

    pub fn lower_cvt(&self, inst: &Value, dst_ty: PtxType, src_ty: PtxType) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let src: u32 = self.get_vreg_for_operand(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::Cvt as u32);
        mi.push_reg(rd);
        mi.push_reg(src);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_cvt_rp(&self, inst: &Value) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let src: u32 = self.get_vreg_for_operand(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::CvtRp as u32);
        mi.push_reg(rd);
        mi.push_reg(src);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_cvt_rn(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::CvtRn as u32)
    }

    pub fn lower_cvt_rz(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::CvtRz as u32)
    }

    pub fn lower_cvt_rm(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::CvtRm as u32)
    }

    pub fn lower_mov(&self, inst: &Value) -> MachineInstr {
        self.lower_two_reg(inst, NvptxOpcode::Mov as u32)
    }

    pub fn lower_cvta(&self, inst: &Value, space: PtxAddrSpace) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let src: u32 = self.get_vreg_for_operand(inst, 0);
        let mut mi = MachineInstr::new(NvptxOpcode::Cvta as u32);
        mi.push_reg(rd);
        mi.push_reg(src);
        mi.push_imm(space as i64); // encode address space
        mi.def = Some(rd);
        mi
    }

    // ========================================================================
    // Bit field extract/insert
    // ========================================================================

    pub fn lower_bfe(&self, inst: &Value) -> MachineInstr {
        // bfe rd, a, b, c: rd = (a >> b) & ((1 << c) - 1)
        let rd: u32 = self.get_or_create_vreg(inst);
        let a: u32 = self.get_vreg_for_operand(inst, 0);
        let b: u32 = self.get_vreg_for_operand(inst, 1);
        let c: u32 = self.get_vreg_for_operand(inst, 2);
        let mut mi = MachineInstr::new(NvptxOpcode::Bfe as u32);
        mi.push_reg(rd);
        mi.push_reg(a);
        mi.push_reg(b);
        mi.push_reg(c);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_bfi(&self, inst: &Value) -> MachineInstr {
        // bfi rd, a, b, c, d: insert bits
        let rd: u32 = self.get_or_create_vreg(inst);
        let a: u32 = self.get_vreg_for_operand(inst, 0);
        let b: u32 = self.get_vreg_for_operand(inst, 1);
        let c: u32 = self.get_vreg_for_operand(inst, 2);
        let d: u32 = self.get_vreg_for_operand(inst, 3);
        let mut mi = MachineInstr::new(NvptxOpcode::Bfi as u32);
        mi.push_reg(rd);
        mi.push_reg(a);
        mi.push_reg(b);
        mi.push_reg(c);
        mi.push_reg(d);
        mi.def = Some(rd);
        mi
    }

    pub fn lower_sad(&self, inst: &Value) -> MachineInstr {
        // sum of absolute differences
        self.lower_three_reg(inst, NvptxOpcode::Sad as u32)
    }

    pub fn lower_selp(&self, inst: &Value) -> MachineInstr {
        // selp: conditional select with predicate
        let rd: u32 = self.get_or_create_vreg(inst);
        let a: u32 = self.get_vreg_for_operand(inst, 0);
        let b: u32 = self.get_vreg_for_operand(inst, 1);
        let p: u32 = self.get_pred_vreg(inst, 2);
        let mut mi = MachineInstr::new(NvptxOpcode::Selp as u32);
        mi.push_reg(rd);
        mi.push_reg(a);
        mi.push_reg(b);
        mi.push_reg(p);
        mi.def = Some(rd);
        mi
    }

    // ========================================================================
    // Helper methods
    // ========================================================================

    fn lower_three_reg(&self, inst: &Value, opcode: u32) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let ra: u32 = self.get_vreg_for_operand(inst, 0);
        let rb: u32 = self.get_vreg_for_operand(inst, 1);
        let mut mi = MachineInstr::new(opcode);
        mi.push_reg(rd);
        mi.push_reg(ra);
        mi.push_reg(rb);
        mi.def = Some(rd);
        mi
    }

    fn lower_four_reg(&self, inst: &Value, opcode: u32) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let ra: u32 = self.get_vreg_for_operand(inst, 0);
        let rb: u32 = self.get_vreg_for_operand(inst, 1);
        let rc: u32 = self.get_vreg_for_operand(inst, 2);
        let mut mi = MachineInstr::new(opcode);
        mi.push_reg(rd);
        mi.push_reg(ra);
        mi.push_reg(rb);
        mi.push_reg(rc);
        mi.def = Some(rd);
        mi
    }

    fn lower_two_reg(&self, inst: &Value, opcode: u32) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let ra: u32 = self.get_vreg_for_operand(inst, 0);
        let mut mi = MachineInstr::new(opcode);
        mi.push_reg(rd);
        mi.push_reg(ra);
        mi.def = Some(rd);
        mi
    }

    fn lower_two_reg_imm(&self, inst: &Value, opcode: u32) -> MachineInstr {
        let rd: u32 = self.get_or_create_vreg(inst);
        let addr: u32 = self.get_vreg_for_operand(inst, 0);
        let mut mi = MachineInstr::new(opcode);
        mi.push_reg(rd);
        mi.push_reg(addr);
        mi.push_imm(0);
        mi.def = Some(rd);
        mi
    }

    fn lower_cmp_two(&self, inst: &Value, opcode: u32) -> MachineInstr {
        let p: u32 = self.get_or_create_vreg_pred(inst, 0);
        let ra: u32 = self.get_vreg_for_operand(inst, 0);
        let rb: u32 = self.get_vreg_for_operand(inst, 1);
        let mut mi = MachineInstr::new(opcode);
        mi.push_reg(p);
        mi.push_reg(ra);
        mi.push_reg(rb);
        mi.def = Some(p);
        mi
    }

    fn get_or_create_vreg(&self, inst: &Value) -> VirtReg {
        *self.vreg_map.get(&(inst.vid as usize)).unwrap_or(&0)
    }

    fn get_vreg_for_operand(&self, inst: &Value, idx: usize) -> VirtReg {
        if idx >= inst.operands.len() {
            return 0;
        }
        let op_ref = &inst.operands[idx];
        let op = op_ref.borrow();
        *self.vreg_map.get(&(op.vid as usize)).unwrap_or(&0)
    }

    fn get_pred_vreg(&self, inst: &Value, idx: usize) -> VirtReg {
        if idx >= inst.operands.len() {
            return 0;
        }
        let op_ref = &inst.operands[idx];
        let op = op_ref.borrow();
        *self.pred_map.get(&(op.vid as usize)).unwrap_or(&0)
    }

    fn get_or_create_vreg_pred(&self, inst: &Value, idx: usize) -> VirtReg {
        *self.pred_map.get(&(inst.vid as usize)).unwrap_or(&0)
    }
}

// ============================================================================
// Compute Capability Pattern Database
// ============================================================================

/// Feature query for conditional lowering based on compute capability
pub struct CcPattern {
    pub required_cc: ComputeCapability,
    pub feature_name: &'static str,
    pub enabled: bool,
}

impl CcPattern {
    pub fn for_cc(cc: ComputeCapability) -> Vec<CcPattern> {
        vec![
            CcPattern {
                required_cc: ComputeCapability::Sm30,
                feature_name: "shuffle",
                enabled: cc >= ComputeCapability::Sm30,
            },
            CcPattern {
                required_cc: ComputeCapability::Sm50,
                feature_name: "funnel_shift",
                enabled: cc >= ComputeCapability::Sm50,
            },
            CcPattern {
                required_cc: ComputeCapability::Sm60,
                feature_name: "dp_fma",
                enabled: cc.supports_dp_fma(),
            },
            CcPattern {
                required_cc: ComputeCapability::Sm70,
                feature_name: "mma",
                enabled: cc.supports_mma(),
            },
            CcPattern {
                required_cc: ComputeCapability::Sm70,
                feature_name: "tensor_core",
                enabled: cc.supports_tensor_core(),
            },
            CcPattern {
                required_cc: ComputeCapability::Sm75,
                feature_name: "ldmatrix",
                enabled: cc.supports_ldmatrix(),
            },
            CcPattern {
                required_cc: ComputeCapability::Sm80,
                feature_name: "sparse_tensor",
                enabled: cc.supports_sparse_tensor_core(),
            },
        ]
    }
}

// ============================================================================
// Tests
// ============================================================================

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

    #[test]
    fn test_compute_capability_from_pair() {
        assert_eq!(
            ComputeCapability::from_pair(7, 0),
            Some(ComputeCapability::Sm70)
        );
        assert_eq!(
            ComputeCapability::from_pair(8, 0),
            Some(ComputeCapability::Sm80)
        );
        assert_eq!(
            ComputeCapability::from_pair(9, 0),
            Some(ComputeCapability::Sm90)
        );
        assert!(ComputeCapability::from_pair(1, 0).is_none());
    }

    #[test]
    fn test_compute_capability_from_name() {
        assert_eq!(
            ComputeCapability::from_name("sm_70"),
            Some(ComputeCapability::Sm70)
        );
        assert_eq!(
            ComputeCapability::from_name("sm_80"),
            Some(ComputeCapability::Sm80)
        );
        assert_eq!(
            ComputeCapability::from_name("sm_90a"),
            Some(ComputeCapability::Sm90a)
        );
    }

    #[test]
    fn test_cc_features() {
        let sm70 = ComputeCapability::Sm70;
        assert!(sm70.supports_mma());
        assert!(sm70.supports_tensor_core());
        assert!(!sm70.supports_sparse_tensor_core());
        assert!(!sm70.supports_ldmatrix());
    }

    #[test]
    fn test_cc_sm80_features() {
        let sm80 = ComputeCapability::Sm80;
        assert!(sm80.supports_sparse_tensor_core());
        assert!(sm80.supports_ldmatrix());
    }

    #[test]
    fn test_cc_warp_size() {
        assert_eq!(ComputeCapability::Sm70.warp_size(), 32);
    }

    #[test]
    fn test_addr_space_mapping() {
        assert_eq!(PtxAddrSpace::from_u32(1), Some(PtxAddrSpace::Global));
        assert_eq!(PtxAddrSpace::from_u32(3), Some(PtxAddrSpace::Shared));
        assert_eq!(PtxAddrSpace::from_u32(4), Some(PtxAddrSpace::Const));
        assert!(PtxAddrSpace::from_u32(99).is_none());
    }

    #[test]
    fn test_ptx_type_sizes() {
        assert_eq!(PtxType::B32.size_bytes(), 4);
        assert_eq!(PtxType::F64.size_bytes(), 8);
        assert_eq!(PtxType::F16.size_bytes(), 2);
    }

    #[test]
    fn test_ptx_type_strings() {
        assert_eq!(PtxType::F32.to_str(), "f32");
        assert_eq!(PtxType::U32.to_str(), "u32");
        assert_eq!(PtxType::S64.to_str(), "s64");
    }

    #[test]
    fn test_atomic_op_strings() {
        assert_eq!(PtxAtomicOp::Add.to_str(), "add");
        assert_eq!(PtxAtomicOp::Cas.to_str(), "cas");
    }

    #[test]
    fn test_selector_creation() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        assert_eq!(sel.cc, ComputeCapability::Sm80);
        assert!(sel.vreg_map.is_empty());
        assert!(!sel.use_fast_math);
    }

    #[test]
    fn test_selector_with_fast_math() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm75).with_fast_math(true);
        assert!(sel.use_fast_math);
    }

    #[test]
    fn test_cc_patterns() {
        let patterns = CcPattern::for_cc(ComputeCapability::Sm80);
        assert!(!patterns.is_empty());
        // All sm80 features should be enabled
        assert!(patterns.iter().all(|p| p.enabled));
    }

    #[test]
    fn test_cc_patterns_sm50() {
        let patterns = CcPattern::for_cc(ComputeCapability::Sm50);
        let mma = patterns.iter().find(|p| p.feature_name == "mma").unwrap();
        assert!(!mma.enabled);
    }

    #[test]
    fn test_shuffle_lowering_stub() {
        use crate::value::Value;
        // Shuffle (shfl) would use vote or similar warp-level instructions
        // This test validates the interface exists
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm70);
        assert_eq!(sel.cc.warp_size(), 32);
    }

    #[test]
    fn test_ldmatrix_support_sm80() {
        let sm80 = ComputeCapability::Sm80;
        assert!(sm80.supports_ldmatrix());
    }

    #[test]
    fn test_tensor_core_support() {
        assert!(ComputeCapability::Sm70.supports_tensor_core());
        assert!(!ComputeCapability::Sm60.supports_tensor_core());
    }

    #[test]
    fn test_num_registers_per_sm() {
        assert_eq!(ComputeCapability::Sm70.num_registers_per_sm(), 65536);
        assert_eq!(ComputeCapability::Sm30.num_registers_per_sm(), 65536);
    }

    #[test]
    fn test_num_warps_per_sm() {
        assert_eq!(ComputeCapability::Sm80.num_warps_per_sm(), 64);
    }

    #[test]
    fn test_all_addr_spaces() {
        assert!(PtxAddrSpace::from_u32(0).is_some());
        assert!(PtxAddrSpace::from_u32(1).is_some());
        assert!(PtxAddrSpace::from_u32(3).is_some());
        assert!(PtxAddrSpace::from_u32(99).is_none());
    }

    #[test]
    fn test_atomic_op_mapping() {
        assert_eq!(PtxAtomicOp::Add.to_str(), "add");
        assert_eq!(PtxAtomicOp::Sub.to_str(), "sub");
        assert_eq!(PtxAtomicOp::Min.to_str(), "min");
        assert_eq!(PtxAtomicOp::Max.to_str(), "max");
        assert_eq!(PtxAtomicOp::And.to_str(), "and");
        assert_eq!(PtxAtomicOp::Or.to_str(), "or");
        assert_eq!(PtxAtomicOp::Xor.to_str(), "xor");
        assert_eq!(PtxAtomicOp::Inc.to_str(), "inc");
        assert_eq!(PtxAtomicOp::Dec.to_str(), "dec");
    }

    #[test]
    fn test_ptx_type_all_strings() {
        assert_eq!(PtxType::B8.to_str(), "b8");
        assert_eq!(PtxType::B16.to_str(), "b16");
        assert_eq!(PtxType::B32.to_str(), "b32");
        assert_eq!(PtxType::B64.to_str(), "b64");
        assert_eq!(PtxType::U8.to_str(), "u8");
        assert_eq!(PtxType::U16.to_str(), "u16");
        assert_eq!(PtxType::U32.to_str(), "u32");
        assert_eq!(PtxType::U64.to_str(), "u64");
        assert_eq!(PtxType::S8.to_str(), "s8");
        assert_eq!(PtxType::S16.to_str(), "s16");
        assert_eq!(PtxType::S32.to_str(), "s32");
        assert_eq!(PtxType::S64.to_str(), "s64");
        assert_eq!(PtxType::F16.to_str(), "f16");
        assert_eq!(PtxType::F32.to_str(), "f32");
        assert_eq!(PtxType::F64.to_str(), "f64");
        assert_eq!(PtxType::Pred.to_str(), "pred");
        assert_eq!(PtxType::B128.to_str(), "b128");
    }

    #[test]
    fn test_addr_space_to_ptx_qualifier() {
        assert_eq!(PtxAddrSpace::Global.to_ptx_qualifier(), ".global");
        assert_eq!(PtxAddrSpace::Shared.to_ptx_qualifier(), ".shared");
        assert_eq!(PtxAddrSpace::Local.to_ptx_qualifier(), ".local");
        assert_eq!(PtxAddrSpace::Const.to_ptx_qualifier(), ".const");
        assert_eq!(PtxAddrSpace::Param.to_ptx_qualifier(), ".param");
    }

    #[test]
    fn test_cc_major_minor() {
        assert_eq!(ComputeCapability::Sm70.major(), 7);
        assert_eq!(ComputeCapability::Sm70.minor(), 0);
        assert_eq!(ComputeCapability::Sm80.major(), 8);
        assert_eq!(ComputeCapability::Sm90.major(), 9);
    }

    #[test]
    fn test_cc_sparse_tensor_sm80() {
        assert!(ComputeCapability::Sm80.supports_sparse_tensor_core());
        assert!(!ComputeCapability::Sm70.supports_sparse_tensor_core());
    }

    #[test]
    fn test_lower_barrier_ops() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        assert_eq!(sel.lower_membar("cta").opcode, NvptxOpcode::Membar as u32);
        assert_eq!(sel.lower_membar("gl").opcode, NvptxOpcode::Membar as u32);
        assert_eq!(sel.lower_membar("sys").opcode, NvptxOpcode::Membar as u32);
    }

    #[test]
    fn test_lower_control_flow_ops() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        assert_eq!(sel.lower_ret().opcode, NvptxOpcode::Ret as u32);
        assert_eq!(sel.lower_exit().opcode, NvptxOpcode::Exit as u32);
        assert_eq!(sel.lower_brkpt().opcode, NvptxOpcode::Brkpt as u32);
    }

    #[test]
    fn test_lower_texture_ops_exist() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        // Just verify the methods exist and return valid opcodes
        let mi = sel.lower_tex(&crate::value::Value::placeholder());
        assert_eq!(mi.opcode, NvptxOpcode::Tex as u32);
    }

    #[test]
    fn test_cc_from_pair_edge_cases() {
        assert!(ComputeCapability::from_pair(0, 0).is_none());
        assert!(ComputeCapability::from_pair(10, 0).is_none());
        assert!(ComputeCapability::from_pair(99, 99).is_none());
    }

    #[test]
    fn test_default_addr_space_is_global() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        assert_eq!(sel.default_addr_space, PtxAddrSpace::Global);
    }

    #[test]
    fn test_fast_math_flag() {
        let sel1 = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        assert!(!sel1.use_fast_math);
        let sel2 = sel1.with_fast_math(true);
        assert!(sel2.use_fast_math);
    }

    #[test]
    fn test_cc_from_invalid_name() {
        assert!(ComputeCapability::from_name("invalid").is_none());
        assert!(ComputeCapability::from_name("sm_99").is_none());
        assert!(ComputeCapability::from_name("").is_none());
    }

    #[test]
    fn test_vote_operations_exist() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        let mi = sel.lower_vote_all(&crate::value::Value::placeholder());
        assert_eq!(mi.opcode, NvptxOpcode::VoteAll as u32);
    }

    #[test]
    fn test_red_operations_exist() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        let mi = sel.lower_red(&crate::value::Value::placeholder(), PtxAtomicOp::Add);
        assert_eq!(mi.opcode, NvptxOpcode::RedAdd as u32);
    }

    #[test]
    fn test_cvt_operations() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        let mi_rp = sel.lower_cvt_rp(&crate::value::Value::placeholder());
        assert_eq!(mi_rp.opcode, NvptxOpcode::CvtRp as u32);
        let mi_rn = sel.lower_cvt_rn(&crate::value::Value::placeholder());
        assert_eq!(mi_rn.opcode, NvptxOpcode::CvtRn as u32);
        let mi_rz = sel.lower_cvt_rz(&crate::value::Value::placeholder());
        assert_eq!(mi_rz.opcode, NvptxOpcode::CvtRz as u32);
    }

    #[test]
    fn test_mov_simple() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        let mi = sel.lower_mov(&crate::value::Value::placeholder());
        assert_eq!(mi.opcode, NvptxOpcode::Mov as u32);
    }

    #[test]
    fn test_cvta_op() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        let mi = sel.lower_cvta(&crate::value::Value::placeholder(), PtxAddrSpace::Shared);
        assert_eq!(mi.opcode, NvptxOpcode::Cvta as u32);
    }

    #[test]
    fn test_atomic_lower_inner() {
        let sel = NvptxFullInstructionSelector::new(ComputeCapability::Sm80);
        let mi = sel.lower_atomic_inner(
            &crate::value::Value::placeholder(),
            PtxAtomicOp::Add,
            PtxAddrSpace::Global,
        );
        assert_eq!(mi.opcode, NvptxOpcode::AtomAdd as u32);
    }
}