llvm-native-core 0.1.16

//! AMDGPU ↔ X86 Bridge Module
//!
//! Clean-room behavioral reconstruction providing a bi-directional bridge
//! between AMDGPU (GCN/RDNA) and X86 ISA semantics. This module supports:
//!
//! - Cross-target instruction mapping (AMDGPU ↔ X86)
//! - Unified register allocation across GPU and CPU register files
//! - Calling convention translation (kernel args → X86 ABI, function call → kernel launch)
//! - Scalar/vector divergence analysis for SIMT-to-SIMD mapping
//! - Shared code generation abstractions for heterogeneous targets
//!
//! Zero LLVM source code consultation. Based on AMD GCN3/RDNA ISA manuals,
//! AMD ROCm documentation, and x86-64 ABI specifications.
//!
//! Architecture coverage:
//! - GCN (Graphics Core Next): GFX6–GFX9
//! - RDNA 1/2/3: GFX1010–GFX1150
//! - Register files: SGPRs (0-105), VGPRs (0-255), AGPRs (0-255)
//! - All 12+ instruction encoding formats
//! - Full calling convention support (kernel + device functions)

use std::collections::{BTreeMap, HashMap, HashSet};
use std::fmt;
use std::sync::Arc;

// ═══════════════════════════════════════════════════════════════════════════════
// Re-export commonly used types from sibling modules
// ═══════════════════════════════════════════════════════════════════════════════

use super::amdgpu_instr_info::{
    AmdgpuEncodingFormat, AmdgpuInstrDesc, AmdgpuInstrInfo, AmdgpuOpcode,
};
use super::amdgpu_register_info::{AmdgpuRegClass, AmdgpuRegisterDesc, AmdgpuRegisterInfo};
use super::amdgpu_target_machine::{
    AmdgpuGeneration, AmdgpuIsaVersion, AmdgpuSubtargetFeatures, AmdgpuTargetMachine,
};

// ═══════════════════════════════════════════════════════════════════════════════
// X86 Register Enumeration (for bridge mapping)
// ═══════════════════════════════════════════════════════════════════════════════

/// X86-64 physical registers for cross-target mapping.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, PartialOrd, Ord)]
pub enum X86PhysReg {
    // General-purpose registers (64-bit)
    RAX,
    RBX,
    RCX,
    RDX,
    RSI,
    RDI,
    RBP,
    RSP,
    R8,
    R9,
    R10,
    R11,
    R12,
    R13,
    R14,
    R15,
    // 32-bit views
    EAX,
    EBX,
    ECX,
    EDX,
    ESI,
    EDI,
    EBP,
    ESP,
    R8D,
    R9D,
    R10D,
    R11D,
    R12D,
    R13D,
    R14D,
    R15D,
    // 16-bit views
    AX,
    BX,
    CX,
    DX,
    SI,
    DI,
    BP,
    SP,
    R8W,
    R9W,
    R10W,
    R11W,
    R12W,
    R13W,
    R14W,
    R15W,
    // 8-bit views
    AL,
    BL,
    CL,
    DL,
    SIL,
    DIL,
    BPL,
    SPL,
    R8B,
    R9B,
    R10B,
    R11B,
    R12B,
    R13B,
    R14B,
    R15B,
    AH,
    BH,
    CH,
    DH,
    // SSE/AVX registers
    XMM0,
    XMM1,
    XMM2,
    XMM3,
    XMM4,
    XMM5,
    XMM6,
    XMM7,
    XMM8,
    XMM9,
    XMM10,
    XMM11,
    XMM12,
    XMM13,
    XMM14,
    XMM15,
    // AVX-512 registers
    YMM0,
    YMM1,
    YMM2,
    YMM3,
    YMM4,
    YMM5,
    YMM6,
    YMM7,
    YMM8,
    YMM9,
    YMM10,
    YMM11,
    YMM12,
    YMM13,
    YMM14,
    YMM15,
    ZMM0,
    ZMM1,
    ZMM2,
    ZMM3,
    ZMM4,
    ZMM5,
    ZMM6,
    ZMM7,
    ZMM8,
    ZMM9,
    ZMM10,
    ZMM11,
    ZMM12,
    ZMM13,
    ZMM14,
    ZMM15,
    ZMM16,
    ZMM17,
    ZMM18,
    ZMM19,
    ZMM20,
    ZMM21,
    ZMM22,
    ZMM23,
    ZMM24,
    ZMM25,
    ZMM26,
    ZMM27,
    ZMM28,
    ZMM29,
    ZMM30,
    ZMM31,
    // x87 FP stack
    ST0,
    ST1,
    ST2,
    ST3,
    ST4,
    ST5,
    ST6,
    ST7,
    // SIMD mask registers (AVX-512)
    K0,
    K1,
    K2,
    K3,
    K4,
    K5,
    K6,
    K7,
    // Flags
    EFLAGS,
}

impl X86PhysReg {
    /// Get the register class for this register.
    pub fn reg_class(&self) -> X86RegClass {
        use X86PhysReg::*;
        match self {
            RAX | RBX | RCX | RDX | RSI | RDI | RBP | RSP | R8 | R9 | R10 | R11 | R12 | R13
            | R14 | R15 | EAX | EBX | ECX | EDX | ESI | EDI | EBP | ESP | R8D | R9D | R10D
            | R11D | R12D | R13D | R14D | R15D | AX | BX | CX | DX | SI | DI | BP | SP | R8W
            | R9W | R10W | R11W | R12W | R13W | R14W | R15W | AL | BL | CL | DL | SIL | DIL
            | BPL | SPL | R8B | R9B | R10B | R11B | R12B | R13B | R14B | R15B | AH | BH | CH
            | DH => X86RegClass::GPR,
            XMM0 | XMM1 | XMM2 | XMM3 | XMM4 | XMM5 | XMM6 | XMM7 | XMM8 | XMM9 | XMM10 | XMM11
            | XMM12 | XMM13 | XMM14 | XMM15 => X86RegClass::XMM,
            YMM0 | YMM1 | YMM2 | YMM3 | YMM4 | YMM5 | YMM6 | YMM7 | YMM8 | YMM9 | YMM10 | YMM11
            | YMM12 | YMM13 | YMM14 | YMM15 => X86RegClass::YMM,
            ZMM0 | ZMM1 | ZMM2 | ZMM3 | ZMM4 | ZMM5 | ZMM6 | ZMM7 | ZMM8 | ZMM9 | ZMM10 | ZMM11
            | ZMM12 | ZMM13 | ZMM14 | ZMM15 | ZMM16 | ZMM17 | ZMM18 | ZMM19 | ZMM20 | ZMM21
            | ZMM22 | ZMM23 | ZMM24 | ZMM25 | ZMM26 | ZMM27 | ZMM28 | ZMM29 | ZMM30 | ZMM31 => {
                X86RegClass::ZMM
            }
            ST0 | ST1 | ST2 | ST3 | ST4 | ST5 | ST6 | ST7 => X86RegClass::X87,
            K0 | K1 | K2 | K3 | K4 | K5 | K6 | K7 => X86RegClass::Mask,
            EFLAGS => X86RegClass::Flags,
        }
    }

    /// Get the size in bits.
    pub fn size_bits(&self) -> u32 {
        use X86PhysReg::*;
        match self {
            RAX | RBX | RCX | RDX | RSI | RDI | RBP | RSP | R8 | R9 | R10 | R11 | R12 | R13
            | R14 | R15 => 64,
            EAX | EBX | ECX | EDX | ESI | EDI | EBP | ESP | R8D | R9D | R10D | R11D | R12D
            | R13D | R14D | R15D => 32,
            AX | BX | CX | DX | SI | DI | BP | SP | R8W | R9W | R10W | R11W | R12W | R13W
            | R14W | R15W => 16,
            AL | BL | CL | DL | SIL | DIL | BPL | SPL | R8B | R9B | R10B | R11B | R12B | R13B
            | R14B | R15B | AH | BH | CH | DH => 8,
            XMM0 | XMM1 | XMM2 | XMM3 | XMM4 | XMM5 | XMM6 | XMM7 | XMM8 | XMM9 | XMM10 | XMM11
            | XMM12 | XMM13 | XMM14 | XMM15 => 128,
            YMM0 | YMM1 | YMM2 | YMM3 | YMM4 | YMM5 | YMM6 | YMM7 | YMM8 | YMM9 | YMM10 | YMM11
            | YMM12 | YMM13 | YMM14 | YMM15 => 256,
            ZMM0 | ZMM1 | ZMM2 | ZMM3 | ZMM4 | ZMM5 | ZMM6 | ZMM7 | ZMM8 | ZMM9 | ZMM10 | ZMM11
            | ZMM12 | ZMM13 | ZMM14 | ZMM15 | ZMM16 | ZMM17 | ZMM18 | ZMM19 | ZMM20 | ZMM21
            | ZMM22 | ZMM23 | ZMM24 | ZMM25 | ZMM26 | ZMM27 | ZMM28 | ZMM29 | ZMM30 | ZMM31 => 512,
            ST0 | ST1 | ST2 | ST3 | ST4 | ST5 | ST6 | ST7 => 80,
            K0 | K1 | K2 | K3 | K4 | K5 | K6 | K7 => 64,
            EFLAGS => 64,
        }
    }

    /// Get the register name string.
    pub fn name(&self) -> &'static str {
        use X86PhysReg::*;
        match self {
            RAX => "rax",
            RBX => "rbx",
            RCX => "rcx",
            RDX => "rdx",
            RSI => "rsi",
            RDI => "rdi",
            RBP => "rbp",
            RSP => "rsp",
            R8 => "r8",
            R9 => "r9",
            R10 => "r10",
            R11 => "r11",
            R12 => "r12",
            R13 => "r13",
            R14 => "r14",
            R15 => "r15",
            EAX => "eax",
            EBX => "ebx",
            ECX => "ecx",
            EDX => "edx",
            ESI => "esi",
            EDI => "edi",
            EBP => "ebp",
            ESP => "esp",
            R8D => "r8d",
            R9D => "r9d",
            R10D => "r10d",
            R11D => "r11d",
            R12D => "r12d",
            R13D => "r13d",
            R14D => "r14d",
            R15D => "r15d",
            AX => "ax",
            BX => "bx",
            CX => "cx",
            DX => "dx",
            SI => "si",
            DI => "di",
            BP => "bp",
            SP => "sp",
            R8W => "r8w",
            R9W => "r9w",
            R10W => "r10w",
            R11W => "r11w",
            R12W => "r12w",
            R13W => "r13w",
            R14W => "r14w",
            R15W => "r15w",
            AL => "al",
            BL => "bl",
            CL => "cl",
            DL => "dl",
            SIL => "sil",
            DIL => "dil",
            BPL => "bpl",
            SPL => "spl",
            R8B => "r8b",
            R9B => "r9b",
            R10B => "r10b",
            R11B => "r11b",
            R12B => "r12b",
            R13B => "r13b",
            R14B => "r14b",
            R15B => "r15b",
            AH => "ah",
            BH => "bh",
            CH => "ch",
            DH => "dh",
            XMM0 => "xmm0",
            XMM1 => "xmm1",
            XMM2 => "xmm2",
            XMM3 => "xmm3",
            XMM4 => "xmm4",
            XMM5 => "xmm5",
            XMM6 => "xmm6",
            XMM7 => "xmm7",
            XMM8 => "xmm8",
            XMM9 => "xmm9",
            XMM10 => "xmm10",
            XMM11 => "xmm11",
            XMM12 => "xmm12",
            XMM13 => "xmm13",
            XMM14 => "xmm14",
            XMM15 => "xmm15",
            YMM0 => "ymm0",
            YMM1 => "ymm1",
            YMM2 => "ymm2",
            YMM3 => "ymm3",
            YMM4 => "ymm4",
            YMM5 => "ymm5",
            YMM6 => "ymm6",
            YMM7 => "ymm7",
            YMM8 => "ymm8",
            YMM9 => "ymm9",
            YMM10 => "ymm10",
            YMM11 => "ymm11",
            YMM12 => "ymm12",
            YMM13 => "ymm13",
            YMM14 => "ymm14",
            YMM15 => "ymm15",
            ZMM0 => "zmm0",
            ZMM1 => "zmm1",
            ZMM2 => "zmm2",
            ZMM3 => "zmm3",
            ZMM4 => "zmm4",
            ZMM5 => "zmm5",
            ZMM6 => "zmm6",
            ZMM7 => "zmm7",
            ZMM8 => "zmm8",
            ZMM9 => "zmm9",
            ZMM10 => "zmm10",
            ZMM11 => "zmm11",
            ZMM12 => "zmm12",
            ZMM13 => "zmm13",
            ZMM14 => "zmm14",
            ZMM15 => "zmm15",
            ZMM16 => "zmm16",
            ZMM17 => "zmm17",
            ZMM18 => "zmm18",
            ZMM19 => "zmm19",
            ZMM20 => "zmm20",
            ZMM21 => "zmm21",
            ZMM22 => "zmm22",
            ZMM23 => "zmm23",
            ZMM24 => "zmm24",
            ZMM25 => "zmm25",
            ZMM26 => "zmm26",
            ZMM27 => "zmm27",
            ZMM28 => "zmm28",
            ZMM29 => "zmm29",
            ZMM30 => "zmm30",
            ZMM31 => "zmm31",
            ST0 => "st(0)",
            ST1 => "st(1)",
            ST2 => "st(2)",
            ST3 => "st(3)",
            ST4 => "st(4)",
            ST5 => "st(5)",
            ST6 => "st(6)",
            ST7 => "st(7)",
            K0 => "k0",
            K1 => "k1",
            K2 => "k2",
            K3 => "k3",
            K4 => "k4",
            K5 => "k5",
            K6 => "k6",
            K7 => "k7",
            EFLAGS => "eflags",
        }
    }

    /// Get the encoding number for this register.
    pub fn encoding(&self) -> u32 {
        use X86PhysReg::*;
        match self {
            RAX | EAX | AX | AL => 0,
            RCX | ECX | CX | CL => 1,
            RDX | EDX | DX | DL => 2,
            RBX | EBX | BX | BL => 3,
            RSP | ESP | SP | SPL => 4,
            RBP | EBP | BP | BPL => 5,
            RSI | ESI | SI | SIL => 6,
            RDI | EDI | DI | DIL => 7,
            R8 | R8D | R8W | R8B => 8,
            R9 | R9D | R9W | R9B => 9,
            R10 | R10D | R10W | R10B => 10,
            R11 | R11D | R11W | R11B => 11,
            R12 | R12D | R12W | R12B => 12,
            R13 | R13D | R13W | R13B => 13,
            R14 | R14D | R14W | R14B => 14,
            R15 | R15D | R15W | R15B => 15,
            AH => 4,
            BH => 5,
            CH => 6,
            DH => 7,
            XMM0 | XMM1 | XMM2 | XMM3 | XMM4 | XMM5 | XMM6 | XMM7 => 0,
            XMM8 | XMM9 | XMM10 | XMM11 | XMM12 | XMM13 | XMM14 | XMM15 => 8,
            YMM0 | YMM1 | YMM2 | YMM3 | YMM4 | YMM5 | YMM6 | YMM7 => 0,
            YMM8 | YMM9 | YMM10 | YMM11 | YMM12 | YMM13 | YMM14 | YMM15 => 8,
            ZMM0 | ZMM1 | ZMM2 | ZMM3 | ZMM4 | ZMM5 | ZMM6 | ZMM7 => 0,
            ZMM8 | ZMM9 | ZMM10 | ZMM11 | ZMM12 | ZMM13 | ZMM14 | ZMM15 => 8,
            ZMM16 | ZMM17 | ZMM18 | ZMM19 | ZMM20 | ZMM21 | ZMM22 | ZMM23 => 16,
            ZMM24 | ZMM25 | ZMM26 | ZMM27 | ZMM28 | ZMM29 | ZMM30 | ZMM31 => 24,
            ST0 => 0,
            ST1 => 1,
            ST2 => 2,
            ST3 => 3,
            ST4 => 4,
            ST5 => 5,
            ST6 => 6,
            ST7 => 7,
            K0 => 0,
            K1 => 1,
            K2 => 2,
            K3 => 3,
            K4 => 4,
            K5 => 5,
            K6 => 6,
            K7 => 7,
            EFLAGS => 0,
            _ => 0,
        }
    }
}

impl fmt::Display for X86PhysReg {
    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
        write!(f, "{}", self.name())
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// X86 Register Class
// ═══════════════════════════════════════════════════════════════════════════════

/// X86-64 register classes.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum X86RegClass {
    /// General-purpose registers (RAX, RBX, etc.)
    GPR,
    /// 128-bit SSE registers
    XMM,
    /// 256-bit AVX registers
    YMM,
    /// 512-bit AVX-512 registers
    ZMM,
    /// x87 floating-point stack
    X87,
    /// AVX-512 mask registers
    Mask,
    /// EFLAGS register
    Flags,
}

impl fmt::Display for X86RegClass {
    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
        match self {
            X86RegClass::GPR => write!(f, "GPR"),
            X86RegClass::XMM => write!(f, "XMM"),
            X86RegClass::YMM => write!(f, "YMM"),
            X86RegClass::ZMM => write!(f, "ZMM"),
            X86RegClass::X87 => write!(f, "X87"),
            X86RegClass::Mask => write!(f, "MASK"),
            X86RegClass::Flags => write!(f, "FLAGS"),
        }
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// X86 Instruction Enumeration
// ═══════════════════════════════════════════════════════════════════════════════

/// X86 instruction opcodes for cross-target mapping.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum X86Opcode {
    // Data movement
    MovRR,
    MovRI,
    MovMR,
    MovRM,
    MovZX,
    MovSX,
    MovSXD,
    Push,
    Pop,
    Xchg,
    Lea,
    Cmov,
    CmovE,
    CmovNE,
    CmovG,
    CmovGE,
    CmovL,
    CmovLE,
    CmovA,
    CmovAE,
    CmovB,
    CmovBE,
    // Arithmetic
    Add,
    Sub,
    Mul,
    IMul,
    Div,
    IDiv,
    Neg,
    Inc,
    Dec,
    // Bitwise
    And,
    Or,
    Xor,
    Not,
    Shl,
    Shr,
    Sar,
    Rol,
    Ror,
    // Comparison
    Cmp,
    Test,
    // Control flow
    Jmp,
    JmpE,
    JmpNE,
    JmpG,
    JmpGE,
    JmpL,
    JmpLE,
    JmpA,
    JmpAE,
    JmpB,
    JmpBE,
    Call,
    Ret,
    // Floating-point (SSE/AVX)
    Addps,
    Subps,
    Mulps,
    Divps,
    Sqrtps,
    Addpd,
    Subpd,
    Mulpd,
    Divpd,
    Sqrtpd,
    Addss,
    Subss,
    Mulss,
    Divss,
    Sqrtss,
    Addsd,
    Subsd,
    Mulsd,
    Divsd,
    Sqrtsd,
    Cvtss2sd,
    Cvtsd2ss,
    Cvtsi2ss,
    Cvtsi2sd,
    Cvttss2si,
    Cvttsd2si,
    Minps,
    Maxps,
    Minss,
    Maxss,
    Andps,
    Andnps,
    Orps,
    Xorps,
    // AVX/AVX2
    VAddps,
    VSubps,
    VMulps,
    VDivps,
    VAddpd,
    VSubpd,
    VMulpd,
    VDivpd,
    VBroadcastss,
    VBroadcastsd,
    VPermilps,
    VPerm2f128,
    VGatherdps,
    VGatherdpd,
    VGatherqps,
    VGatherqpd,
    // AVX-512
    VAddpsZ,
    VSubpsZ,
    VMulpsZ,
    VDivpsZ,
    VAddpdZ,
    VSubpdZ,
    VMulpdZ,
    VDivpdZ,
    // Memory
    Nop,
}

impl X86Opcode {
    pub fn mnemonic(&self) -> &'static str {
        use X86Opcode::*;
        match self {
            MovRR => "mov",
            MovRI => "mov",
            MovMR => "mov",
            MovRM => "mov",
            MovZX => "movzx",
            MovSX => "movsx",
            MovSXD => "movsxd",
            Push => "push",
            Pop => "pop",
            Xchg => "xchg",
            Lea => "lea",
            Cmov => "cmov",
            CmovE => "cmove",
            CmovNE => "cmovne",
            CmovG => "cmovg",
            CmovGE => "cmovge",
            CmovL => "cmovl",
            CmovLE => "cmovle",
            CmovA => "cmova",
            CmovAE => "cmovae",
            CmovB => "cmovb",
            CmovBE => "cmovbe",
            Add => "add",
            Sub => "sub",
            Mul => "mul",
            IMul => "imul",
            Div => "div",
            IDiv => "idiv",
            Neg => "neg",
            Inc => "inc",
            Dec => "dec",
            And => "and",
            Or => "or",
            Xor => "xor",
            Not => "not",
            Shl => "shl",
            Shr => "shr",
            Sar => "sar",
            Rol => "rol",
            Ror => "ror",
            Cmp => "cmp",
            Test => "test",
            Jmp => "jmp",
            JmpE => "je",
            JmpNE => "jne",
            JmpG => "jg",
            JmpGE => "jge",
            JmpL => "jl",
            JmpLE => "jle",
            JmpA => "ja",
            JmpAE => "jae",
            JmpB => "jb",
            JmpBE => "jbe",
            Call => "call",
            Ret => "ret",
            Addps => "addps",
            Subps => "subps",
            Mulps => "mulps",
            Divps => "divps",
            Sqrtps => "sqrtps",
            Addpd => "addpd",
            Subpd => "subpd",
            Mulpd => "mulpd",
            Divpd => "divpd",
            Sqrtpd => "sqrtpd",
            Addss => "addss",
            Subss => "subss",
            Mulss => "mulss",
            Divss => "divss",
            Sqrtss => "sqrtss",
            Addsd => "addsd",
            Subsd => "subsd",
            Mulsd => "mulsd",
            Divsd => "divsd",
            Sqrtsd => "sqrtsd",
            Cvtss2sd => "cvtss2sd",
            Cvtsd2ss => "cvtsd2ss",
            Cvtsi2ss => "cvtsi2ss",
            Cvtsi2sd => "cvtsi2sd",
            Cvttss2si => "cvttss2si",
            Cvttsd2si => "cvttsd2si",
            Minps => "minps",
            Maxps => "maxps",
            Minss => "minss",
            Maxss => "maxss",
            Andps => "andps",
            Andnps => "andnps",
            Orps => "orps",
            Xorps => "xorps",
            VAddps => "vaddps",
            VSubps => "vsubps",
            VMulps => "vmulps",
            VDivps => "vdivps",
            VAddpd => "vaddpd",
            VSubpd => "vsubpd",
            VMulpd => "vmulpd",
            VDivpd => "vdivpd",
            VBroadcastss => "vbroadcastss",
            VBroadcastsd => "vbroadcastsd",
            VPermilps => "vpermilps",
            VPerm2f128 => "vperm2f128",
            VGatherdps => "vgatherdps",
            VGatherdpd => "vgatherdpd",
            VGatherqps => "vgatherqps",
            VGatherqpd => "vgatherqpd",
            VAddpsZ => "vaddps",
            VSubpsZ => "vsubps",
            VMulpsZ => "vmulps",
            VDivpsZ => "vdivps",
            VAddpdZ => "vaddpd",
            VSubpdZ => "vsubpd",
            VMulpdZ => "vmulpd",
            VDivpdZ => "vdivpd",
            Nop => "nop",
        }
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// Register Mapping: AMDGPU Register <-> X86 Register
// ═══════════════════════════════════════════════════════════════════════════════

/// Maps an AMDGPU register to its corresponding X86 register for cross-target translation.
/// SGPRs map to GPRs, VGPRs map to XMM/YMM/ZMM registers.
#[derive(Debug, Clone)]
pub struct RegisterMapping {
    pub amdgpu_reg_id: u32,
    pub amdgpu_reg_name: String,
    pub amdgpu_reg_class: AmdgpuRegClass,
    pub x86_reg: X86PhysReg,
    pub is_direct: bool,
    pub requires_swizzle: bool,
}

impl RegisterMapping {
    pub fn new(
        amdgpu_reg_id: u32,
        amdgpu_reg_name: &str,
        amdgpu_reg_class: AmdgpuRegClass,
        x86_reg: X86PhysReg,
    ) -> Self {
        Self {
            amdgpu_reg_id,
            amdgpu_reg_name: amdgpu_reg_name.to_string(),
            amdgpu_reg_class,
            x86_reg,
            is_direct: true,
            requires_swizzle: false,
        }
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// Instruction Mapping: AMDGPU Instruction <-> X86 Instruction(s)
// ═══════════════════════════════════════════════════════════════════════════════

/// Describes how an AMDGPU instruction maps to X86 instruction(s).
#[derive(Debug, Clone)]
pub struct InstructionMapping {
    pub amdgpu_opcode: AmdgpuOpcode,
    pub x86_opcodes: Vec<X86Opcode>,
    /// Number of X86 instructions needed to emulate this AMDGPU instruction.
    pub expansion_factor: u32,
    /// Whether SIMT execution affects the mapping (needs mask/lane handling).
    pub simt_aware: bool,
    /// Whether this maps 1:1 or requires multiple X86 instructions.
    pub is_direct: bool,
    /// Whether the mapping preserves exact semantics.
    pub preserves_semantics: bool,
    /// Additional notes about the mapping.
    pub notes: &'static str,
}

impl InstructionMapping {
    pub fn new(
        amdgpu_opcode: AmdgpuOpcode,
        x86_opcodes: Vec<X86Opcode>,
        expansion_factor: u32,
    ) -> Self {
        Self {
            amdgpu_opcode,
            x86_opcodes,
            expansion_factor,
            simt_aware: false,
            is_direct: expansion_factor == 1,
            preserves_semantics: true,
            notes: "",
        }
    }

    pub fn with_simt(mut self, simt: bool) -> Self {
        self.simt_aware = simt;
        self
    }

    pub fn with_notes(mut self, notes: &'static str) -> Self {
        self.notes = notes;
        self
    }

    pub fn with_semantics(mut self, preserves: bool) -> Self {
        self.preserves_semantics = preserves;
        self
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Calling Convention (for cross-target)
// ═══════════════════════════════════════════════════════════════════════════════

/// AMDGPU calling convention types.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum AmdgpuCallingConvention {
    /// Kernel entry point (launched from host).
    Kernel,
    /// Device function (called from other GPU functions).
    DeviceFunction,
    /// Indirect function call.
    Indirect,
    /// Internal convention (for cross-target optimization).
    Fast,
}

impl AmdgpuCallingConvention {
    pub fn name(&self) -> &'static str {
        match self {
            Self::Kernel => "amdgpu_kernel",
            Self::DeviceFunction => "amdgpu_device",
            Self::Indirect => "amdgpu_indirect",
            Self::Fast => "amdgpu_fast",
        }
    }

    /// Whether this CC uses SGPRs for arguments.
    pub fn uses_sgpr_args(&self) -> bool {
        match self {
            Self::Kernel => true,
            Self::DeviceFunction => false,
            Self::Indirect => false,
            Self::Fast => true,
        }
    }

    /// Whether this CC returns via VGPRs.
    pub fn returns_via_vgpr(&self) -> bool {
        match self {
            Self::Kernel => false,
            Self::DeviceFunction => true,
            Self::Indirect => true,
            Self::Fast => true,
        }
    }

    /// Whether 'sret' (struct return) is supported.
    pub fn supports_sret(&self) -> bool {
        match self {
            Self::Kernel => false,
            Self::DeviceFunction => true,
            Self::Indirect => true,
            Self::Fast => false,
        }
    }

    /// Number of SGPR registers used for implicit kernel args.
    pub fn implicit_sgpr_count(&self) -> u32 {
        match self {
            Self::Kernel => 6, // workgroup_id_{x,y,z}, workitem_id_{x,y,z}
            _ => 0,
        }
    }

    /// Number of VGPR registers used for implicit args.
    pub fn implicit_vgpr_count(&self) -> u32 {
        match self {
            Self::Kernel => 1, // workitem_id in VGPR0
            _ => 0,
        }
    }
}

impl fmt::Display for AmdgpuCallingConvention {
    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
        write!(f, "{}", self.name())
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Kernel ABI Argument Descriptor
// ═══════════════════════════════════════════════════════════════════════════════

/// Describes a kernel argument for cross-target ABI translation.
#[derive(Debug, Clone)]
pub struct AmdgpuKernelArg {
    pub name: String,
    pub size_bytes: u32,
    pub alignment: u32,
    pub offset: u32,
    pub is_pointer: bool,
    pub is_byval: bool,
    pub addr_space: AmdgpuAddrSpace,
    pub sgpr_index: Option<u32>,
    pub vgpr_index: Option<u32>,
}

impl AmdgpuKernelArg {
    pub fn new(name: &str, size_bytes: u32, alignment: u32, offset: u32) -> Self {
        Self {
            name: name.to_string(),
            size_bytes,
            alignment,
            offset,
            is_pointer: false,
            is_byval: false,
            addr_space: AmdgpuAddrSpace::Global,
            sgpr_index: None,
            vgpr_index: None,
        }
    }

    pub fn with_pointer(mut self, is_ptr: bool) -> Self {
        self.is_pointer = is_ptr;
        self
    }

    pub fn with_addr_space(mut self, space: AmdgpuAddrSpace) -> Self {
        self.addr_space = space;
        self
    }

    pub fn with_sgpr(mut self, idx: u32) -> Self {
        self.sgpr_index = Some(idx);
        self
    }

    pub fn with_vgpr(mut self, idx: u32) -> Self {
        self.vgpr_index = Some(idx);
        self
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Address Spaces
// ═══════════════════════════════════════════════════════════════════════════════

/// AMDGPU address spaces.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum AmdgpuAddrSpace {
    /// Global memory (VRAM)
    Global = 1,
    /// Constant memory
    Constant = 4,
    /// Local Data Share (LDS / shared memory)
    Local = 3,
    /// Private / scratch memory
    Private = 5,
    /// Generic / flat address space
    Generic = 0,
    /// Region (GDS)
    Region = 2,
}

impl AmdgpuAddrSpace {
    pub fn from_u32(n: u32) -> Option<Self> {
        match n {
            0 => Some(Self::Generic),
            1 => Some(Self::Global),
            2 => Some(Self::Region),
            3 => Some(Self::Local),
            4 => Some(Self::Constant),
            5 => Some(Self::Private),
            _ => None,
        }
    }

    pub fn name(&self) -> &'static str {
        match self {
            Self::Global => "global",
            Self::Constant => "constant",
            Self::Local => "local",
            Self::Private => "private",
            Self::Generic => "generic",
            Self::Region => "region",
        }
    }

    pub fn to_llvm_addrspace(&self) -> u32 {
        *self as u32
    }
}

impl fmt::Display for AmdgpuAddrSpace {
    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
        write!(f, "{}", self.name())
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Register Info (Cross-Target Extended)
// ═══════════════════════════════════════════════════════════════════════════════

/// Extended AMDGPU register information for cross-target operations.
/// Provides register allocation, mapping, and constraint information
/// needed by the X86 bridge.
#[derive(Debug, Clone)]
pub struct AmdgpuRegisterInfoX86 {
    /// Base AMDGPU register info.
    pub base: AmdgpuRegisterInfo,
    /// Mapping from AMDGPU register ID to X86 register.
    pub register_mappings: Vec<RegisterMapping>,
    /// Mapping from X86 register to AMDGPU register ID.
    pub x86_to_amdgpu: HashMap<X86PhysReg, u32>,
    /// Available SGPR count for the current ISA.
    pub available_sgprs: u32,
    /// Available VGPR count for the current ISA.
    pub available_vgprs: u32,
    /// Whether AGPRs are available.
    pub has_agprs: bool,
}

impl AmdgpuRegisterInfoX86 {
    pub fn new(isa: AmdgpuIsaVersion) -> Self {
        let base = AmdgpuRegisterInfo::new();
        let available_sgprs = isa.max_sgprs();
        let available_vgprs = isa.max_vgprs();
        let has_agprs = isa.has_agprs();

        // Build register mappings: SGPRs -> GPRs, VGPRs -> XMMs
        let mut mappings = Vec::new();
        let mut x86_to_amdgpu = HashMap::new();

        // Map SGPR0-SGPR15 -> RAX, RCX, RDX, RBX, RSP, RBP, RSI, RDI, R8-R15
        let sgpr_to_gpr: [(u32, X86PhysReg); 16] = [
            (0, X86PhysReg::RAX),
            (1, X86PhysReg::RCX),
            (2, X86PhysReg::RDX),
            (3, X86PhysReg::RBX),
            (4, X86PhysReg::RSP),
            (5, X86PhysReg::RBP),
            (6, X86PhysReg::RSI),
            (7, X86PhysReg::RDI),
            (8, X86PhysReg::R8),
            (9, X86PhysReg::R9),
            (10, X86PhysReg::R10),
            (11, X86PhysReg::R11),
            (12, X86PhysReg::R12),
            (13, X86PhysReg::R13),
            (14, X86PhysReg::R14),
            (15, X86PhysReg::R15),
        ];
        for (sgpr_id, x86_reg) in &sgpr_to_gpr {
            let name = format!("sgpr{}", sgpr_id);
            let mapping = RegisterMapping::new(*sgpr_id, &name, AmdgpuRegClass::SGPR32, *x86_reg);
            x86_to_amdgpu.insert(*x86_reg, *sgpr_id);
            mappings.push(mapping);
        }

        // Map VGPR0-VGPR15 -> XMM0-XMM15
        for i in 0..16u32 {
            let xmm_reg = match i {
                0 => X86PhysReg::XMM0,
                1 => X86PhysReg::XMM1,
                2 => X86PhysReg::XMM2,
                3 => X86PhysReg::XMM3,
                4 => X86PhysReg::XMM4,
                5 => X86PhysReg::XMM5,
                6 => X86PhysReg::XMM6,
                7 => X86PhysReg::XMM7,
                8 => X86PhysReg::XMM8,
                9 => X86PhysReg::XMM9,
                10 => X86PhysReg::XMM10,
                11 => X86PhysReg::XMM11,
                12 => X86PhysReg::XMM12,
                13 => X86PhysReg::XMM13,
                14 => X86PhysReg::XMM14,
                15 => X86PhysReg::XMM15,
                _ => unreachable!(),
            };
            let name = format!("vgpr{}", i);
            let mapping = RegisterMapping::new(i + 256, &name, AmdgpuRegClass::VGPR32, xmm_reg);
            x86_to_amdgpu.insert(xmm_reg, i + 256);
            mappings.push(mapping);
        }

        Self {
            base,
            register_mappings: mappings,
            x86_to_amdgpu,
            available_sgprs,
            available_vgprs,
            has_agprs,
        }
    }

    /// Get the X86 register for a given AMDGPU register.
    pub fn map_to_x86(&self, amdgpu_reg_id: u32) -> Option<X86PhysReg> {
        self.register_mappings
            .iter()
            .find(|m| m.amdgpu_reg_id == amdgpu_reg_id)
            .map(|m| m.x86_reg)
    }

    /// Get the AMDGPU register for a given X86 register.
    pub fn map_from_x86(&self, x86_reg: X86PhysReg) -> Option<u32> {
        self.x86_to_amdgpu.get(&x86_reg).copied()
    }

    /// Get the total number of mapped registers.
    pub fn mapped_count(&self) -> usize {
        self.register_mappings.len()
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Cross-Target Target Machine
// ═══════════════════════════════════════════════════════════════════════════════

/// Extended AMDGPU target machine with X86 bridge capabilities.
/// Wraps the base target machine and adds cross-target translation support.
#[derive(Debug, Clone)]
pub struct AmdgpuX86TargetMachine {
    /// The base AMDGPU target machine.
    pub base: AmdgpuTargetMachine,
    /// X86 register information for mapping.
    pub x86_register_info: AmdgpuRegisterInfoX86,
    /// Instruction mappings from AMDGPU to X86.
    pub instruction_mappings: Vec<InstructionMapping>,
    /// The target triple for X86 code generation.
    pub x86_triple: String,
    /// The X86 data layout string.
    pub x86_data_layout: String,
    /// Optimization level for cross-target code generation.
    pub opt_level: u32,
    /// Whether to use AVX instructions.
    pub use_avx: bool,
    /// Whether to use AVX2 instructions.
    pub use_avx2: bool,
    /// Whether to use AVX-512 instructions.
    pub use_avx512: bool,
    /// Whether to emit X86 code with SIMT lane masks.
    pub emit_simt_masks: bool,
}

impl AmdgpuX86TargetMachine {
    pub fn new(isa: AmdgpuIsaVersion) -> Self {
        let base = AmdgpuTargetMachine::new(isa);
        let x86_register_info = AmdgpuRegisterInfoX86::new(isa);
        let instruction_mappings = build_default_instruction_mappings();

        Self {
            base,
            x86_register_info,
            instruction_mappings,
            x86_triple: "x86_64-unknown-linux-gnu".to_string(),
            x86_data_layout:
                "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
                    .to_string(),
            opt_level: 2,
            use_avx: true,
            use_avx2: true,
            use_avx512: false,
            emit_simt_masks: true,
        }
    }

    pub fn from_cpu(cpu: &str) -> Option<Self> {
        AmdgpuIsaVersion::from_str(cpu).map(Self::new)
    }

    /// Set the X86 optimization level.
    pub fn set_opt_level(&mut self, level: u32) {
        self.opt_level = level.min(3);
    }

    /// Enable AVX-512 for wider vector mapping.
    pub fn enable_avx512(&mut self) {
        self.use_avx512 = true;
    }

    /// Get the best XMM/YMM/ZMM register for a VGPR based on available width.
    pub fn map_vgpr_to_x86_vector(&self, vgpr_id: u32) -> Option<X86PhysReg> {
        let xmm = self.x86_register_info.map_to_x86(vgpr_id)?;
        // Upgrade based on available features
        if self.use_avx512 {
            // Map to ZMM if AVX-512 is available
            let offset = match xmm {
                X86PhysReg::XMM0 => Some(X86PhysReg::ZMM0),
                X86PhysReg::XMM1 => Some(X86PhysReg::ZMM1),
                X86PhysReg::XMM2 => Some(X86PhysReg::ZMM2),
                X86PhysReg::XMM3 => Some(X86PhysReg::ZMM3),
                X86PhysReg::XMM4 => Some(X86PhysReg::ZMM4),
                X86PhysReg::XMM5 => Some(X86PhysReg::ZMM5),
                X86PhysReg::XMM6 => Some(X86PhysReg::ZMM6),
                X86PhysReg::XMM7 => Some(X86PhysReg::ZMM7),
                X86PhysReg::XMM8 => Some(X86PhysReg::ZMM8),
                X86PhysReg::XMM9 => Some(X86PhysReg::ZMM9),
                X86PhysReg::XMM10 => Some(X86PhysReg::ZMM10),
                X86PhysReg::XMM11 => Some(X86PhysReg::ZMM11),
                X86PhysReg::XMM12 => Some(X86PhysReg::ZMM12),
                X86PhysReg::XMM13 => Some(X86PhysReg::ZMM13),
                X86PhysReg::XMM14 => Some(X86PhysReg::ZMM14),
                X86PhysReg::XMM15 => Some(X86PhysReg::ZMM15),
                _ => None,
            };
            return offset;
        } else if self.use_avx2 || self.use_avx {
            let offset = match xmm {
                X86PhysReg::XMM0 => Some(X86PhysReg::YMM0),
                X86PhysReg::XMM1 => Some(X86PhysReg::YMM1),
                X86PhysReg::XMM2 => Some(X86PhysReg::YMM2),
                X86PhysReg::XMM3 => Some(X86PhysReg::YMM3),
                X86PhysReg::XMM4 => Some(X86PhysReg::YMM4),
                X86PhysReg::XMM5 => Some(X86PhysReg::YMM5),
                X86PhysReg::XMM6 => Some(X86PhysReg::YMM6),
                X86PhysReg::XMM7 => Some(X86PhysReg::YMM7),
                X86PhysReg::XMM8 => Some(X86PhysReg::YMM8),
                X86PhysReg::XMM9 => Some(X86PhysReg::YMM9),
                X86PhysReg::XMM10 => Some(X86PhysReg::YMM10),
                X86PhysReg::XMM11 => Some(X86PhysReg::YMM11),
                X86PhysReg::XMM12 => Some(X86PhysReg::YMM12),
                X86PhysReg::XMM13 => Some(X86PhysReg::YMM13),
                X86PhysReg::XMM14 => Some(X86PhysReg::YMM14),
                X86PhysReg::XMM15 => Some(X86PhysReg::YMM15),
                _ => None,
            };
            return offset;
        }
        Some(xmm)
    }

    /// Get the feature string for the X86 target.
    pub fn x86_feature_string(&self) -> String {
        let mut features = Vec::new();
        features.push("+sse2".to_string());
        if self.use_avx {
            features.push("+avx".to_string());
        }
        if self.use_avx2 {
            features.push("+avx2".to_string());
        }
        if self.use_avx512 {
            features.push("+avx512f".to_string());
            features.push("+avx512dq".to_string());
            features.push("+avx512bw".to_string());
            features.push("+avx512vl".to_string());
        }
        features.join(",")
    }

    /// Get the data layout for X86.
    pub fn get_x86_data_layout(&self) -> &str {
        &self.x86_data_layout
    }

    /// Look up the instruction mapping for a given AMDGPU opcode.
    pub fn lookup_mapping(&self, opcode: AmdgpuOpcode) -> Option<&InstructionMapping> {
        self.instruction_mappings
            .iter()
            .find(|m| m.amdgpu_opcode == opcode)
    }

    /// Get all instruction mappings.
    pub fn all_mappings(&self) -> &[InstructionMapping] {
        &self.instruction_mappings
    }

    /// Count the number of instruction mappings.
    pub fn mapping_count(&self) -> usize {
        self.instruction_mappings.len()
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// Default Instruction Mappings: AMDGPU -> X86
// ═══════════════════════════════════════════════════════════════════════════════

fn build_default_instruction_mappings() -> Vec<InstructionMapping> {
    let mut mappings = Vec::new();

    // ── SOP1 scalar instructions ───────────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SMovB32,
        vec![X86Opcode::MovRR],
        1,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::SMovB64,
            vec![X86Opcode::MovRR, X86Opcode::MovRR],
            2,
        )
        .with_notes("64-bit move requires two 32-bit moves"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SNotB32,
        vec![X86Opcode::Not],
        1,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::SAbsI32,
            vec![X86Opcode::MovRR, X86Opcode::Neg, X86Opcode::CmovL],
            3,
        )
        .with_notes("abs = (x < 0) ? -x : x"),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SBrevB32, vec![X86Opcode::MovRR], 1)
            .with_notes("bit-reverse -> use BSWAP + ROR pattern"),
    );

    // ── SOP2 scalar ALU instructions ───────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SAddU32,
        vec![X86Opcode::Add],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SSubU32,
        vec![X86Opcode::Sub],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SAddI32,
        vec![X86Opcode::Add],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SSubI32,
        vec![X86Opcode::Sub],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SAndB32,
        vec![X86Opcode::And],
        1,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::SAndB64,
            vec![X86Opcode::And, X86Opcode::And],
            2,
        )
        .with_notes("64-bit AND via two 32-bit"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SOrB32,
        vec![X86Opcode::Or],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SOrB64,
        vec![X86Opcode::Or, X86Opcode::Or],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SXorB32,
        vec![X86Opcode::Xor],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SXorB64,
        vec![X86Opcode::Xor, X86Opcode::Xor],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SLshlB32,
        vec![X86Opcode::Shl],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SLshrB32,
        vec![X86Opcode::Shr],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SAshrI32,
        vec![X86Opcode::Sar],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SMulI32,
        vec![X86Opcode::IMul],
        1,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::SMinU32,
            vec![X86Opcode::Cmp, X86Opcode::CmovB],
            2,
        )
        .with_notes("minu via cmp+cmovb"),
    );
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::SMinI32,
            vec![X86Opcode::Cmp, X86Opcode::CmovL],
            2,
        )
        .with_notes("min via cmp+cmovl"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SMaxU32,
        vec![X86Opcode::Cmp, X86Opcode::CmovA],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SMaxI32,
        vec![X86Opcode::Cmp, X86Opcode::CmovG],
        2,
    ));
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCselectB32, vec![X86Opcode::CmovNE], 1)
            .with_notes("cselect via cmov"),
    );

    // ── SOPC comparison instructions ───────────────────────────────────────
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpEqI32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpLgI32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpGtI32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpGeI32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpLtI32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpLeI32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpEqU32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpGtU32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCmpLtU32, vec![X86Opcode::Cmp], 1).with_simt(true),
    );

    // ── SOPK instructions ──────────────────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SMovkI32,
        vec![X86Opcode::MovRI],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SAddkI32,
        vec![X86Opcode::Add],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SMulkI32,
        vec![X86Opcode::IMul],
        1,
    ));

    // ── SOPP instructions ──────────────────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SNop,
        vec![X86Opcode::Nop],
        1,
    ));
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SBranch, vec![X86Opcode::Jmp], 1).with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCbranchScc0, vec![X86Opcode::JmpE], 1)
            .with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCbranchScc1, vec![X86Opcode::JmpNE], 1)
            .with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCbranchVccz, vec![X86Opcode::JmpE], 1)
            .with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCbranchVccnz, vec![X86Opcode::JmpNE], 1)
            .with_simt(true),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCbranchExecz, vec![X86Opcode::JmpE], 1)
            .with_simt(true)
            .with_notes("EXEC=0 means all lanes masked → jmp if mask==0"),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::SCbranchExecnz, vec![X86Opcode::JmpNE], 1)
            .with_simt(true),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SEndpgm,
        vec![X86Opcode::Ret],
        1,
    ));

    // ── VOP1 vector instructions ───────────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VNop,
        vec![X86Opcode::Nop],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMovB32,
        vec![X86Opcode::MovRR],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMovB64,
        vec![X86Opcode::MovRR, X86Opcode::MovRR],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCvtF32I32,
        vec![X86Opcode::Cvtsi2ss],
        1,
    ));
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::VCvtF32U32, vec![X86Opcode::Cvtsi2ss], 1)
            .with_notes("unsigned conversion needs adjustment for large values"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCvtI32F32,
        vec![X86Opcode::Cvttss2si],
        1,
    ));
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::VCvtU32F32, vec![X86Opcode::Cvttss2si], 1)
            .with_notes("unsigned conversion via cvttss2si + fixup"),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::VCvtF16F32, vec![X86Opcode::Cvtss2sd], 1)
            .with_semantics(false),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::VCvtF32F16, vec![X86Opcode::Cvtsd2ss], 1)
            .with_semantics(false),
    );
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VRcpF32,
            vec![X86Opcode::MovRI, X86Opcode::Divss],
            2,
        )
        .with_notes("rcp = 1.0 / x; approximate on GPU vs precise on CPU"),
    );
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VRsqF32,
            vec![X86Opcode::Sqrtss, X86Opcode::MovRI, X86Opcode::Divss],
            3,
        )
        .with_notes("rsq = 1/sqrt(x)"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VSqrtF32,
        vec![X86Opcode::Sqrtss],
        1,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VCeilF32,
            vec![
                X86Opcode::MovRR,
                X86Opcode::Cvttss2si,
                X86Opcode::Cvtsi2ss,
                X86Opcode::Addss,
            ],
            4,
        )
        .with_notes("ceil via software"),
    );
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VFloorF32,
            vec![
                X86Opcode::MovRR,
                X86Opcode::Cvttss2si,
                X86Opcode::Cvtsi2ss,
                X86Opcode::Subss,
            ],
            4,
        )
        .with_notes("floor via software"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VTruncF32,
        vec![X86Opcode::Cvttss2si, X86Opcode::Cvtsi2ss],
        2,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VRndneF32,
            vec![X86Opcode::Cvtss2si, X86Opcode::Cvtsi2ss],
            2,
        )
        .with_notes("round-to-nearest-even"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VFractF32,
        vec![X86Opcode::Cvttss2si, X86Opcode::Cvtsi2ss, X86Opcode::Subss],
        3,
    ));

    // ── VOP2 vector ALU instructions ───────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VAddF32,
        vec![X86Opcode::Addss],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VSubF32,
        vec![X86Opcode::Subss],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMulF32,
        vec![X86Opcode::Mulss],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMinF32,
        vec![X86Opcode::Minss],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMaxF32,
        vec![X86Opcode::Maxss],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VAndB32,
        vec![X86Opcode::And],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VOrB32,
        vec![X86Opcode::Or],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VXorB32,
        vec![X86Opcode::Xor],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VLshlB32,
        vec![X86Opcode::Shl],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VLshrB32,
        vec![X86Opcode::Shr],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VAshrI32,
        vec![X86Opcode::Sar],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VLshlrevB32,
        vec![X86Opcode::Shl],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VLshrrevB32,
        vec![X86Opcode::Shr],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VAshrrevI32,
        vec![X86Opcode::Sar],
        1,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VAddcU32,
            vec![X86Opcode::Add, X86Opcode::Add],
            2,
        )
        .with_notes("add with carry"),
    );
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VSubbU32,
            vec![X86Opcode::Sub, X86Opcode::Sub],
            2,
        )
        .with_notes("sub with borrow"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMacF32,
        vec![X86Opcode::Mulss, X86Opcode::Addss],
        2,
    ));

    // ── VOPC vector comparison instructions ────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpF32Eq,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpF32Lg,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpF32Gt,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpF32Ge,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpF32Lt,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpF32Le,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpU32Eq,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpU32Gt,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpU32Lt,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpI32Eq,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpI32Gt,
        vec![X86Opcode::Cmp],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VCmpI32Lt,
        vec![X86Opcode::Cmp],
        1,
    ));

    // ── VOP3 extended vector instructions ──────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VAddF64,
        vec![X86Opcode::Addsd],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VSubF64,
        vec![X86Opcode::Subsd],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMulF64,
        vec![X86Opcode::Mulsd],
        1,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VFmaF32,
            vec![X86Opcode::Mulss, X86Opcode::Addss],
            2,
        )
        .with_notes("FMA = a*b + c; x86 without FMA expands to mul+add"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VFmaF64,
        vec![X86Opcode::Mulsd, X86Opcode::Addsd],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMadF32,
        vec![X86Opcode::Mulss, X86Opcode::Addss],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VMadF64,
        vec![X86Opcode::Mulsd, X86Opcode::Addsd],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::VBfeU32,
        vec![X86Opcode::Shr, X86Opcode::And],
        2,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VBfiB32,
            vec![
                X86Opcode::And,
                X86Opcode::Shl,
                X86Opcode::And,
                X86Opcode::Or,
            ],
            4,
        )
        .with_notes("bitfield insert"),
    );

    // ── VOP3P packed instructions ──────────────────────────────────────────
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::VPkAddF16, vec![X86Opcode::Addps], 1)
            .with_semantics(false)
            .with_notes("packed f16 add: use f32 lanes for each half on x86"),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::VPkSubF16, vec![X86Opcode::Subps], 1)
            .with_semantics(false),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::VPkMulF16, vec![X86Opcode::Mulps], 1)
            .with_semantics(false),
    );
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::VPkFmaF16,
            vec![X86Opcode::Mulps, X86Opcode::Addps],
            2,
        )
        .with_semantics(false),
    );

    // ── SMEM scalar memory instructions ────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SLoadDword,
        vec![X86Opcode::MovRM],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SLoadDwordX2,
        vec![X86Opcode::MovRM, X86Opcode::MovRM],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SLoadDwordX4,
        vec![
            X86Opcode::MovRM,
            X86Opcode::MovRM,
            X86Opcode::MovRM,
            X86Opcode::MovRM,
        ],
        4,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SStoreDword,
        vec![X86Opcode::MovMR],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SStoreDwordX2,
        vec![X86Opcode::MovMR, X86Opcode::MovMR],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::SStoreDwordX4,
        vec![
            X86Opcode::MovMR,
            X86Opcode::MovMR,
            X86Opcode::MovMR,
            X86Opcode::MovMR,
        ],
        4,
    ));

    // ── FLAT memory instructions ───────────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatLoadDword,
        vec![X86Opcode::MovRM],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatLoadDwordX2,
        vec![X86Opcode::MovRM, X86Opcode::MovRM],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatLoadDwordX4,
        vec![
            X86Opcode::MovRM,
            X86Opcode::MovRM,
            X86Opcode::MovRM,
            X86Opcode::MovRM,
        ],
        4,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatStoreDword,
        vec![X86Opcode::MovMR],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatStoreDwordX2,
        vec![X86Opcode::MovMR, X86Opcode::MovMR],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatStoreDwordX4,
        vec![
            X86Opcode::MovMR,
            X86Opcode::MovMR,
            X86Opcode::MovMR,
            X86Opcode::MovMR,
        ],
        4,
    ));
    mappings.push(
        InstructionMapping::new(
            AmdgpuOpcode::FlatAtomicAdd,
            vec![X86Opcode::MovRM, X86Opcode::Add, X86Opcode::MovMR],
            3,
        )
        .with_notes("atomic via lock cmpxchg loop"),
    );
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatAtomicSub,
        vec![X86Opcode::MovRM, X86Opcode::Sub, X86Opcode::MovMR],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatAtomicAnd,
        vec![X86Opcode::MovRM, X86Opcode::And, X86Opcode::MovMR],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatAtomicOr,
        vec![X86Opcode::MovRM, X86Opcode::Or, X86Opcode::MovMR],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::FlatAtomicXor,
        vec![X86Opcode::MovRM, X86Opcode::Xor, X86Opcode::MovMR],
        3,
    ));

    // ── DS (LDS) instructions ──────────────────────────────────────────────
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsWriteB32,
        vec![X86Opcode::MovMR],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsWriteB64,
        vec![X86Opcode::MovMR, X86Opcode::MovMR],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsReadB32,
        vec![X86Opcode::MovRM],
        1,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsReadB64,
        vec![X86Opcode::MovRM, X86Opcode::MovRM],
        2,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsAddU32,
        vec![X86Opcode::MovRM, X86Opcode::Add, X86Opcode::MovMR],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsSubU32,
        vec![X86Opcode::MovRM, X86Opcode::Sub, X86Opcode::MovMR],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsMinU32,
        vec![X86Opcode::MovRM, X86Opcode::Cmp, X86Opcode::CmovB],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsMaxU32,
        vec![X86Opcode::MovRM, X86Opcode::Cmp, X86Opcode::CmovA],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsAndB32,
        vec![X86Opcode::MovRM, X86Opcode::And, X86Opcode::MovMR],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsOrB32,
        vec![X86Opcode::MovRM, X86Opcode::Or, X86Opcode::MovMR],
        3,
    ));
    mappings.push(InstructionMapping::new(
        AmdgpuOpcode::DsXorB32,
        vec![X86Opcode::MovRM, X86Opcode::Xor, X86Opcode::MovMR],
        3,
    ));

    // ── MIMG (image) instructions ──────────────────────────────────────────
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::ImageLoad, vec![X86Opcode::MovRM], 1)
            .with_notes("image load expands to texture sampling calls on x86"),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::ImageStore, vec![X86Opcode::MovMR], 1)
            .with_notes("image store expands to texture write calls on x86"),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::ImageSample, vec![X86Opcode::Call], 1)
            .with_semantics(false)
            .with_notes("image sample requires full texture sampling call"),
    );

    // ── EXP instructions ───────────────────────────────────────────────────
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::Exp, vec![X86Opcode::Call], 1)
            .with_semantics(false)
            .with_notes("export maps to render target write or function call"),
    );
    mappings.push(
        InstructionMapping::new(AmdgpuOpcode::ExpPos0, vec![X86Opcode::Call], 1)
            .with_semantics(false),
    );

    mappings
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU ↔ X86 Bridge — Main Bridge Struct
// ═══════════════════════════════════════════════════════════════════════════════

/// The main AMDGPU ↔ X86 bridge, providing bidirectional translation
/// between AMDGPU and X86 code generation.
#[derive(Debug, Clone)]
pub struct AMDGPUX86Bridge {
    /// The AMDGPU target machine with X86 bridge extensions.
    pub target_machine: AmdgpuX86TargetMachine,
    /// The current calling convention.
    pub calling_convention: AmdgpuCallingConvention,
    /// The AMDGPU instruction info.
    pub amdgpu_instr_info: AmdgpuInstrInfo,
    /// Kernel arguments for the current function.
    pub kernel_args: Vec<AmdgpuKernelArg>,
    /// Cross-target codegen context.
    pub cross_target: CrossTargetAmdgpu,
    /// Statistics for bridge operations.
    pub stats: BridgeStats,
}

impl AMDGPUX86Bridge {
    /// Create a new bridge for the given ISA version.
    pub fn new(isa: AmdgpuIsaVersion) -> Self {
        let target_machine = AmdgpuX86TargetMachine::new(isa);
        let amdgpu_instr_info = AmdgpuInstrInfo::new();
        let cross_target = CrossTargetAmdgpu::new(isa);

        Self {
            target_machine,
            calling_convention: AmdgpuCallingConvention::Kernel,
            amdgpu_instr_info,
            kernel_args: Vec::new(),
            cross_target,
            stats: BridgeStats::default(),
        }
    }

    /// Create a bridge from a CPU name string.
    pub fn from_cpu(cpu: &str) -> Option<Self> {
        AmdgpuIsaVersion::from_str(cpu).map(Self::new)
    }

    /// Set the calling convention for the bridge.
    pub fn set_calling_convention(&mut self, cc: AmdgpuCallingConvention) {
        self.calling_convention = cc;
    }

    /// Add a kernel argument to the bridge.
    pub fn add_kernel_arg(&mut self, arg: AmdgpuKernelArg) {
        self.kernel_args.push(arg);
    }

    /// Clear all kernel arguments.
    pub fn clear_kernel_args(&mut self) {
        self.kernel_args.clear();
    }

    /// Translate an AMDGPU opcode to its X86 equivalent(s).
    /// Returns the X86 opcode(s) and the expansion factor.
    pub fn translate_instruction(
        &self,
        amdgpu_opcode: AmdgpuOpcode,
    ) -> Option<(Vec<X86Opcode>, u32)> {
        self.target_machine
            .lookup_mapping(amdgpu_opcode)
            .map(|m| (m.x86_opcodes.clone(), m.expansion_factor))
    }

    /// Check if an AMDGPU instruction has a direct X86 mapping.
    pub fn has_direct_mapping(&self, opcode: AmdgpuOpcode) -> bool {
        self.target_machine
            .lookup_mapping(opcode)
            .map(|m| m.is_direct)
            .unwrap_or(false)
    }

    /// Get the total expansion factor for a sequence of AMDGPU instructions.
    pub fn compute_expansion(&self, opcodes: &[AmdgpuOpcode]) -> u32 {
        opcodes
            .iter()
            .filter_map(|op| self.target_machine.lookup_mapping(*op))
            .map(|m| m.expansion_factor)
            .sum()
    }

    /// Get the ISA version of the bridge.
    pub fn isa_version(&self) -> AmdgpuIsaVersion {
        self.target_machine.base.isa_version
    }

    /// Whether the bridge supports wave32 mode.
    pub fn supports_wave32(&self) -> bool {
        self.isa_version().supports_wave32()
    }

    /// Enable AVX-512 features.
    pub fn enable_avx512(&mut self) {
        self.target_machine.enable_avx512();
        self.stats.avx512_enabled = true;
    }

    /// Set the optimization level.
    pub fn set_opt_level(&mut self, level: u32) {
        self.target_machine.set_opt_level(level);
    }

    /// Generate the X86 feature string based on current configuration.
    pub fn x86_feature_string(&self) -> String {
        self.target_machine.x86_feature_string()
    }

    /// Get the X86 data layout string.
    pub fn get_x86_data_layout(&self) -> &str {
        self.target_machine.get_x86_data_layout()
    }

    /// Get the AMDGPU data layout (via base target machine).
    pub fn get_amdgpu_triple(&self) -> &str {
        &self.target_machine.base.triple
    }

    /// Emit a prologue for the X86 equivalent of an AMDGPU kernel.
    pub fn emit_x86_prologue(&self) -> String {
        let mut buf = String::new();
        buf.push_str("; -- AMDGPU Kernel Prologue (X86 Bridge) --\n");
        buf.push_str(&format!("; AMDGPU ISA: {}\n", self.isa_version().as_str()));
        buf.push_str(&format!(
            "; X86 Target: {}\n",
            self.target_machine.x86_triple
        ));
        buf.push_str(&format!(
            "; Calling Convention: {}\n",
            self.calling_convention
        ));

        if self.calling_convention == AmdgpuCallingConvention::Kernel {
            buf.push_str("; Kernel Launch Parameters:\n");
            for arg in &self.kernel_args {
                buf.push_str(&format!(
                    ";   arg {} @ offset {} (size {}, align {})\n",
                    arg.name, arg.offset, arg.size_bytes, arg.alignment
                ));
            }
            buf.push_str("; Implicit kernel args in SGPRs: workgroup_id_{x,y,z}\n");
        }

        buf
    }

    /// Emit an epilogue for the X86 equivalent of an AMDGPU kernel.
    pub fn emit_x86_epilogue(&self) -> String {
        let mut buf = String::new();
        buf.push_str("; -- AMDGPU Kernel Epilogue (X86 Bridge) --\n");
        if self.calling_convention == AmdgpuCallingConvention::Kernel {
            buf.push_str("; s_endpgm → ret\n");
        }
        buf
    }

    /// Compute a suggested X86 register allocation for the given VGPR count.
    pub fn suggest_x86_register_allocation(&self, vgpr_count: u32) -> Vec<X86PhysReg> {
        let mut regs = Vec::new();
        // Trivially map VGPRs to XMM registers
        for i in 0..vgpr_count.min(16) {
            if let Some(reg) = self.target_machine.map_vgpr_to_x86_vector(i) {
                regs.push(reg);
            }
        }
        regs
    }

    /// Analyze scalar/vector divergence for a sequence of opcodes.
    pub fn analyze_divergence(&self, opcodes: &[AmdgpuOpcode]) -> DivergenceAnalysis {
        let mut analysis = DivergenceAnalysis::default();

        for &op in opcodes {
            let is_scalar = self
                .amdgpu_instr_info
                .get(op)
                .map(|d| {
                    // Determined by encoding format
                    matches!(
                        d.encoding_format,
                        AmdgpuEncodingFormat::SOP1
                            | AmdgpuEncodingFormat::SOP2
                            | AmdgpuEncodingFormat::SOPC
                            | AmdgpuEncodingFormat::SOPK
                            | AmdgpuEncodingFormat::SOPP
                            | AmdgpuEncodingFormat::SMEM
                    )
                })
                .unwrap_or(false);

            if is_scalar {
                analysis.scalar_count += 1;
            } else {
                analysis.vector_count += 1;
            }
        }

        analysis
    }

    /// Check if the bridge is ready (has valid configuration).
    pub fn is_ready(&self) -> bool {
        !self.target_machine.base.triple.is_empty() && self.target_machine.mapping_count() > 0
    }

    /// Get bridge statistics.
    pub fn get_stats(&self) -> &BridgeStats {
        &self.stats
    }

    /// Reset bridge statistics.
    pub fn reset_stats(&mut self) {
        self.stats = BridgeStats::default();
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// Bridge Statistics
// ═══════════════════════════════════════════════════════════════════════════════

/// Statistics for bridge operations.
#[derive(Debug, Clone, Default)]
pub struct BridgeStats {
    /// Number of instructions translated.
    pub instructions_translated: u64,
    /// Number of kernel calls processed.
    pub kernels_processed: u64,
    /// Total expansion factor across all translations.
    pub total_expansion: u64,
    /// Number of SIMT-aware translations.
    pub simt_translations: u64,
    /// Whether AVX-512 is enabled.
    pub avx512_enabled: bool,
    /// Number of register mappings used.
    pub register_mappings_used: u64,
}

// ═══════════════════════════════════════════════════════════════════════════════
// Divergence Analysis
// ═══════════════════════════════════════════════════════════════════════════════

/// Result of scalar/vector divergence analysis.
#[derive(Debug, Clone, Default)]
pub struct DivergenceAnalysis {
    /// Number of scalar (uniform) operations.
    pub scalar_count: u64,
    /// Number of vector (divergent) operations.
    pub vector_count: u64,
    /// Total operations analyzed.
    pub total_count: u64,
}

impl DivergenceAnalysis {
    pub fn total(&self) -> u64 {
        self.scalar_count + self.vector_count
    }

    pub fn scalar_ratio(&self) -> f64 {
        let total = self.total();
        if total == 0 {
            0.0
        } else {
            self.scalar_count as f64 / total as f64
        }
    }

    pub fn vector_ratio(&self) -> f64 {
        let total = self.total();
        if total == 0 {
            0.0
        } else {
            self.vector_count as f64 / total as f64
        }
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// Cross-Target AMDGPU — Shared Codegen (Scalar/Vector Divergence)
// ═══════════════════════════════════════════════════════════════════════════════

/// Cross-target code generation context for AMDGPU.
/// Manages the distinction between scalar (uniform across wavefront)
/// and vector (per-lane) operations when targeting X86.
#[derive(Debug, Clone)]
pub struct CrossTargetAmdgpu {
    /// The ISA version for this context.
    pub isa_version: AmdgpuIsaVersion,
    /// Whether to scalarize uniform operations.
    pub scalarize_uniform: bool,
    /// Whether to emit SIMT lane masks.
    pub emit_lane_masks: bool,
    /// Wavefront size for SIMT emulation.
    pub wavefront_size: u32,
    /// Maximum VGPRs available.
    pub max_vgprs: u32,
    /// Maximum SGPRs available.
    pub max_sgprs: u32,
    /// Currently mapped VGPRs.
    pub allocated_vgprs: HashSet<u32>,
    /// Currently mapped SGPRs.
    pub allocated_sgprs: HashSet<u32>,
    /// Lane ID for SIMT emulation.
    pub lane_id_reg: Option<X86PhysReg>,
    /// Exec mask for SIMT emulation.
    pub exec_mask_reg: Option<X86PhysReg>,
}

impl CrossTargetAmdgpu {
    pub fn new(isa: AmdgpuIsaVersion) -> Self {
        Self {
            isa_version: isa,
            scalarize_uniform: true,
            emit_lane_masks: true,
            wavefront_size: if isa.supports_wave32() { 32 } else { 64 },
            max_vgprs: isa.max_vgprs(),
            max_sgprs: isa.max_sgprs(),
            allocated_vgprs: HashSet::new(),
            allocated_sgprs: HashSet::new(),
            lane_id_reg: None,
            exec_mask_reg: None,
        }
    }

    /// Allocate a VGPR for cross-target use.
    pub fn allocate_vgpr(&mut self) -> Option<u32> {
        for i in 0..self.max_vgprs {
            if !self.allocated_vgprs.contains(&i) {
                self.allocated_vgprs.insert(i);
                return Some(i);
            }
        }
        None
    }

    /// Allocate an SGPR for cross-target use.
    pub fn allocate_sgpr(&mut self) -> Option<u32> {
        for i in 0..self.max_sgprs {
            if !self.allocated_sgprs.contains(&i) {
                self.allocated_sgprs.insert(i);
                return Some(i);
            }
        }
        None
    }

    /// Free a previously allocated VGPR.
    pub fn free_vgpr(&mut self, reg: u32) {
        self.allocated_vgprs.remove(&reg);
    }

    /// Free a previously allocated SGPR.
    pub fn free_sgpr(&mut self, reg: u32) {
        self.allocated_sgprs.remove(&reg);
    }

    /// Set the lane ID register for SIMT emulation.
    pub fn set_lane_id_reg(&mut self, reg: X86PhysReg) {
        self.lane_id_reg = Some(reg);
    }

    /// Set the exec mask register for SIMT emulation.
    pub fn set_exec_mask_reg(&mut self, reg: X86PhysReg) {
        self.exec_mask_reg = Some(reg);
    }

    /// Determine if an AMDGPU instruction encoding format is scalar.
    pub fn is_scalar_format(format: AmdgpuEncodingFormat) -> bool {
        matches!(
            format,
            AmdgpuEncodingFormat::SOP1
                | AmdgpuEncodingFormat::SOP2
                | AmdgpuEncodingFormat::SOPC
                | AmdgpuEncodingFormat::SOPK
                | AmdgpuEncodingFormat::SOPP
                | AmdgpuEncodingFormat::SMEM
        )
    }

    /// Determine if an AMDGPU instruction encoding format is vector.
    pub fn is_vector_format(format: AmdgpuEncodingFormat) -> bool {
        matches!(
            format,
            AmdgpuEncodingFormat::VOP1
                | AmdgpuEncodingFormat::VOP2
                | AmdgpuEncodingFormat::VOPC
                | AmdgpuEncodingFormat::VOP3
                | AmdgpuEncodingFormat::VOP3P
                | AmdgpuEncodingFormat::FLAT
                | AmdgpuEncodingFormat::DS
                | AmdgpuEncodingFormat::MIMG
                | AmdgpuEncodingFormat::EXP
                | AmdgpuEncodingFormat::VINTRP
        )
    }

    /// Get the number of lanes per wavefront.
    pub fn lanes_per_wavefront(&self) -> u32 {
        self.wavefront_size
    }

    /// Emit SIMT lane mask initialization code for X86.
    pub fn emit_lane_init(&self) -> String {
        if !self.emit_lane_masks {
            return String::new();
        }

        let mut buf = String::new();
        buf.push_str("; -- SIMT Lane Mask Initialization --\n");
        buf.push_str(&format!("; Wavefront size: {}\n", self.wavefront_size));
        buf.push_str("; mov eax, [thread_idx]\n");
        buf.push_str("; and eax, lane_mask\n");
        buf.push_str("; This assigns each logical lane a unique XMM register slot\n");
        buf
    }

    /// Emit exec mask predicate for conditional execution.
    pub fn emit_exec_predicate(&self, inverted: bool) -> String {
        if !self.emit_lane_masks {
            return String::new();
        }

        if inverted {
            "; test exec_mask_reg, exec_mask_reg ; jz skip_block\n".to_string()
        } else {
            "; test exec_mask_reg, exec_mask_reg ; jnz exec_block\n".to_string()
        }
    }

    /// Get the total count of allocated registers.
    pub fn allocated_count(&self) -> usize {
        self.allocated_sgprs.len() + self.allocated_vgprs.len()
    }

    /// Reset all allocations.
    pub fn reset_allocations(&mut self) {
        self.allocated_vgprs.clear();
        self.allocated_sgprs.clear();
        self.lane_id_reg = None;
        self.exec_mask_reg = None;
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Kernel Descriptor (Cross-Target ABI)
// ═══════════════════════════════════════════════════════════════════════════════

/// A full kernel descriptor for ABI translation to X86.
#[derive(Debug, Clone)]
pub struct AmdgpuKernelDescriptorX86 {
    /// Kernel name.
    pub name: String,
    /// Kernel arguments.
    pub args: Vec<AmdgpuKernelArg>,
    /// SGPR usage count.
    pub sgpr_count: u32,
    /// VGPR usage count.
    pub vgpr_count: u32,
    /// SGPR spill size in bytes.
    pub sgpr_spill_size: u32,
    /// VGPR spill size in bytes.
    pub vgpr_spill_size: u32,
    /// Private segment size.
    pub private_segment_size: u32,
    /// Group segment size (LDS).
    pub group_segment_size: u32,
    /// Whether the kernel uses dynamic LDS.
    pub uses_dynamic_lds: bool,
    /// Kernel code entry byte offset.
    pub kernel_code_entry_byte_offset: u64,
    /// Workgroup size hint (X, Y, Z).
    pub workgroup_size_hint: (u32, u32, u32),
    /// Required workgroup size (X, Y, Z), if fixed.
    pub required_workgroup_size: Option<(u32, u32, u32)>,
}

impl AmdgpuKernelDescriptorX86 {
    pub fn new(name: &str) -> Self {
        Self {
            name: name.to_string(),
            args: Vec::new(),
            sgpr_count: 0,
            vgpr_count: 0,
            sgpr_spill_size: 0,
            vgpr_spill_size: 0,
            private_segment_size: 0,
            group_segment_size: 0,
            uses_dynamic_lds: false,
            kernel_code_entry_byte_offset: 256,
            workgroup_size_hint: (64, 1, 1),
            required_workgroup_size: None,
        }
    }

    /// Add a kernel argument.
    pub fn add_arg(&mut self, arg: AmdgpuKernelArg) {
        self.args.push(arg);
    }

    /// Set register usage.
    pub fn set_register_usage(&mut self, sgprs: u32, vgprs: u32) {
        self.sgpr_count = sgprs;
        self.vgpr_count = vgprs;
    }

    /// Set spill sizes.
    pub fn set_spill(&mut self, sgpr_spill: u32, vgpr_spill: u32) {
        self.sgpr_spill_size = sgpr_spill;
        self.vgpr_spill_size = vgpr_spill;
    }

    /// Generate a C-compatible X86 kernel launch stub.
    pub fn emit_x86_launch_stub(&self) -> String {
        let mut buf = String::new();
        buf.push_str(&format!(
            "// X86 Launch Stub for AMDGPU Kernel: {}\n",
            self.name
        ));
        buf.push_str("// This stub would be called by the host-side runtime\n");
        buf.push_str(&format!(
            "// void launch_{}(/* kernel args */);\n",
            self.name
        ));
        buf.push_str("//\n");
        buf.push_str("// amd_hostcall_execute_kernel(\n");
        buf.push_str(&format!("//     \"{}\",\n", self.name));
        buf.push_str(&format!(
            "//     grid_dim, block_dim, shared_mem={},\n",
            self.group_segment_size
        ));
        buf.push_str(&format!(
            "//     sgpr_count={}, vgpr_count={}\n",
            self.sgpr_count, self.vgpr_count
        ));
        buf.push_str("// );\n");
        buf
    }

    /// Convert to a JSON-like string for metadata.
    pub fn to_metadata_string(&self) -> String {
        let mut buf = String::new();
        buf.push_str("{");
        buf.push_str(&format!("\"name\":\"{}\"", self.name));
        buf.push_str(&format!(",\"sgpr_count\":{}", self.sgpr_count));
        buf.push_str(&format!(",\"vgpr_count\":{}", self.vgpr_count));
        buf.push_str(&format!(",\"sgpr_spill\":{}", self.sgpr_spill_size));
        buf.push_str(&format!(",\"vgpr_spill\":{}", self.vgpr_spill_size));
        buf.push_str(&format!(
            ",\"group_segment_size\":{}",
            self.group_segment_size
        ));
        buf.push_str(&format!(
            ",\"private_segment_size\":{}",
            self.private_segment_size
        ));
        buf.push_str(&format!(",\"uses_dynamic_lds\":{}", self.uses_dynamic_lds));
        buf.push_str(&format!(
            ",\"kernel_code_offset\":{}",
            self.kernel_code_entry_byte_offset
        ));
        buf.push_str("}");
        buf
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Kernel ABI Lowering Context (Cross-Target)
// ═══════════════════════════════════════════════════════════════════════════════

/// Context for lowering AMDGPU kernel ABI to X86.
#[derive(Debug, Clone)]
pub struct AmdgpuKernelAbiContext {
    /// The kernel name.
    pub kernel_name: String,
    /// The calling convention.
    pub calling_convention: AmdgpuCallingConvention,
    /// Work-item ID in X (local id within workgroup).
    pub local_id_x: Option<X86PhysReg>,
    /// Work-item ID in Y.
    pub local_id_y: Option<X86PhysReg>,
    /// Work-item ID in Z.
    pub local_id_z: Option<X86PhysReg>,
    /// Workgroup ID in X.
    pub group_id_x: Option<X86PhysReg>,
    /// Workgroup ID in Y.
    pub group_id_y: Option<X86PhysReg>,
    /// Workgroup ID in Z.
    pub group_id_z: Option<X86PhysReg>,
    /// Dispatch thread ID in X (global).
    pub global_id_x: Option<X86PhysReg>,
    /// Dispatch thread ID in Y.
    pub global_id_y: Option<X86PhysReg>,
    /// Dispatch thread ID in Z.
    pub global_id_z: Option<X86PhysReg>,
}

impl AmdgpuKernelAbiContext {
    pub fn new(kernel_name: &str) -> Self {
        Self {
            kernel_name: kernel_name.to_string(),
            calling_convention: AmdgpuCallingConvention::Kernel,
            local_id_x: None,
            local_id_y: None,
            local_id_z: None,
            group_id_x: None,
            group_id_y: None,
            group_id_z: None,
            global_id_x: None,
            global_id_y: None,
            global_id_z: None,
        }
    }

    /// Compute global ID from local ID + group ID + group size.
    /// global_id = group_id * local_size + local_id
    pub fn emit_compute_global_id_x(&self) -> String {
        let mut buf = String::new();
        buf.push_str("; Compute global_id.x = group_id.x * local_size.x + local_id.x\n");
        buf.push_str("; mov eax, [group_id_x]\n");
        buf.push_str("; imul eax, [local_size_x]\n");
        buf.push_str("; add eax, [local_id_x]\n");
        buf
    }

    /// Emit work-item ID loading for X86 (emulating kernel launch).
    pub fn emit_load_workitem_id_x86(&self) -> String {
        let mut buf = String::new();
        buf.push_str("; -- AMDGPU Work-Item ID Loading (X86 Bridge) --\n");
        buf.push_str("; On X86, work-item IDs are thread-local values\n");
        buf.push_str("; assigned by the kernel launch runtime\n");
        if self.calling_convention == AmdgpuCallingConvention::Kernel {
            buf.push_str("; local_id_x is in VGPR0 (first VGPR argument)\n");
            buf.push_str("; local_id_y is in VGPR1\n");
            buf.push_str("; local_id_z is in VGPR2\n");
            buf.push_str("; group_id_x is in SGPR4\n");
            buf.push_str("; group_id_y is in SGPR5\n");
            buf.push_str("; group_id_z is in SGPR6\n");
        }
        buf
    }

    /// Emit dispatch packet descriptor for AMD AQL (HSA AQL).
    pub fn emit_aql_dispatch(&self) -> String {
        let mut buf = String::new();
        buf.push_str("; AMD AQL Dispatch Packet:\n");
        buf.push_str("; struct amd_kernel_dispatch {\n");
        buf.push_str(";     uint16_t header;\n");
        buf.push_str(";     uint16_t setup;\n");
        buf.push_str(";     uint16_t workgroup_x;\n");
        buf.push_str(";     uint16_t workgroup_y;\n");
        buf.push_str(";     uint16_t workgroup_z;\n");
        buf.push_str(";     uint16_t reserved;\n");
        buf.push_str(";     uint32_t grid_size_x;\n");
        buf.push_str(";     uint32_t grid_size_y;\n");
        buf.push_str(";     uint32_t grid_size_z;\n");
        buf.push_str(";     uint32_t private_segment_size;\n");
        buf.push_str(";     uint32_t group_segment_size;\n");
        buf.push_str(";     uint64_t kernel_object;\n");
        buf.push_str(";     uint64_t kernarg_address;\n");
        buf.push_str(";     uint64_t reserved2;\n");
        buf.push_str(";     uint64_t completion_signal;\n");
        buf.push_str("; };\n");
        buf
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// AMDGPU Memory Model Translation
// ═══════════════════════════════════════════════════════════════════════════════

/// Memory ordering for cross-target translation.
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum AmdgpuMemoryOrder {
    Relaxed,
    Acquire,
    Release,
    AcquireRelease,
    SequentiallyConsistent,
}

impl AmdgpuMemoryOrder {
    pub fn to_x86_mfence(&self) -> bool {
        // On x86, only sequentially consistent needs explicit MFENCE
        // (x86-TSO provides acquire/release for free with locked instructions)
        matches!(self, AmdgpuMemoryOrder::SequentiallyConsistent)
    }

    pub fn to_x86_lfence(&self) -> bool {
        // LFENCE may be needed for acquire ordering on some implementations
        matches!(
            self,
            AmdgpuMemoryOrder::Acquire | AmdgpuMemoryOrder::AcquireRelease
        )
    }

    pub fn to_x86_sfence(&self) -> bool {
        // SFENCE for release ordering
        matches!(
            self,
            AmdgpuMemoryOrder::Release | AmdgpuMemoryOrder::AcquireRelease
        )
    }
}

/// Memory scope for AMDGPU operations.
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum AmdgpuMemoryScope {
    /// Single work-item.
    WorkItem,
    /// Single wavefront.
    Wavefront,
    /// Single workgroup.
    Workgroup,
    /// All workgroups on the device.
    Device,
    /// System-wide (includes host CPU).
    System,
}

impl AmdgpuMemoryScope {
    pub fn requires_barrier(&self) -> bool {
        matches!(
            self,
            AmdgpuMemoryScope::Workgroup | AmdgpuMemoryScope::Device | AmdgpuMemoryScope::System
        )
    }

    pub fn to_x86_fence(&self) -> Option<&'static str> {
        match self {
            AmdgpuMemoryScope::WorkItem | AmdgpuMemoryScope::Wavefront => None,
            AmdgpuMemoryScope::Workgroup => Some("__syncthreads()"),
            AmdgpuMemoryScope::Device => Some("mfence"),
            AmdgpuMemoryScope::System => Some("mfence"),
        }
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// SIMT Wavefront Emulation Utilities
// ═══════════════════════════════════════════════════════════════════════════════

/// Utility for emulating AMDGPU wavefront execution on X86 (SIMT→SIMD).
#[derive(Debug, Clone)]
pub struct WavefrontEmulator {
    /// Wavefront size (32 or 64).
    pub wave_size: u32,
    /// Number of waves being emulated.
    pub num_waves: u32,
    /// Current lane ID being processed.
    pub current_lane: u32,
    /// Exec mask (active lanes).
    pub exec_mask: u64,
    /// Whether to scalarize uniform branches.
    pub scalarize_uniform: bool,
}

impl WavefrontEmulator {
    pub fn new(wave_size: u32) -> Self {
        Self {
            wave_size,
            num_waves: 1,
            current_lane: 0,
            exec_mask: if wave_size == 64 {
                u64::MAX
            } else {
                u32::MAX as u64
            },
            scalarize_uniform: true,
        }
    }

    /// Set the exec mask (active lane bitmask).
    pub fn set_exec_mask(&mut self, mask: u64) {
        self.exec_mask = mask;
    }

    /// Check if a given lane is active.
    pub fn lane_active(&self, lane: u32) -> bool {
        (self.exec_mask >> lane) & 1 == 1
    }

    /// Get the number of active lanes.
    pub fn active_lane_count(&self) -> u32 {
        self.exec_mask.count_ones()
    }

    /// Check if all lanes are active.
    pub fn all_lanes_active(&self) -> bool {
        let full_mask = if self.wave_size == 64 {
            u64::MAX
        } else {
            u32::MAX as u64
        };
        self.exec_mask == full_mask
    }

    /// Emit X86 code for a SIMT loop header (iterates over lanes).
    pub fn emit_simt_loop_header(&self) -> String {
        let mut buf = String::new();
        buf.push_str(&format!(
            "; SIMT Wavefront Loop: {} lanes, mask=0x{:016x}\n",
            self.wave_size, self.exec_mask
        ));
        buf.push_str("; .L_simt_loop:\n");
        buf.push_str(";   mov eax, [lane_counter]\n");
        buf.push_str(";   cmp eax, wavefront_size\n");
        buf.push_str(";   jge .L_simt_done\n");
        buf.push_str(";   bt [exec_mask], eax\n");
        buf.push_str(";   jnc .L_lane_inactive\n");
        buf.push_str(";   ; -- active lane body --\n");
        buf
    }

    /// Emit X86 code for a SIMT loop footer.
    pub fn emit_simt_loop_footer(&self) -> String {
        let mut buf = String::new();
        buf.push_str("; .L_lane_inactive:\n");
        buf.push_str(";   inc [lane_counter]\n");
        buf.push_str(";   jmp .L_simt_loop\n");
        buf.push_str("; .L_simt_done:\n");
        buf
    }

    /// Compute the mask for a given subset of lanes.
    pub fn lane_mask(start: u32, count: u32) -> u64 {
        if count == 0 || start >= 64 {
            return 0;
        }
        ((1u64 << count) - 1) << start
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// Complete AMDGPU Bridge Builder
// ═══════════════════════════════════════════════════════════════════════════════

/// Builder for constructing an [`AMDGPUX86Bridge`] with fluent API.
#[derive(Debug, Clone)]
pub struct AmdgpuBridgeBuilder {
    isa: AmdgpuIsaVersion,
    opt_level: u32,
    use_avx: bool,
    use_avx2: bool,
    use_avx512: bool,
    calling_convention: AmdgpuCallingConvention,
    emit_simt: bool,
    scalarize: bool,
    kernel_args: Vec<AmdgpuKernelArg>,
}

impl AmdgpuBridgeBuilder {
    pub fn new(isa: AmdgpuIsaVersion) -> Self {
        Self {
            isa,
            opt_level: 2,
            use_avx: true,
            use_avx2: true,
            use_avx512: false,
            calling_convention: AmdgpuCallingConvention::Kernel,
            emit_simt: true,
            scalarize: true,
            kernel_args: Vec::new(),
        }
    }

    pub fn from_cpu(cpu: &str) -> Option<Self> {
        AmdgpuIsaVersion::from_str(cpu).map(Self::new)
    }

    pub fn opt_level(mut self, level: u32) -> Self {
        self.opt_level = level.min(3);
        self
    }

    pub fn avx(mut self, enable: bool) -> Self {
        self.use_avx = enable;
        self
    }

    pub fn avx2(mut self, enable: bool) -> Self {
        self.use_avx2 = enable;
        self
    }

    pub fn avx512(mut self, enable: bool) -> Self {
        self.use_avx512 = enable;
        self
    }

    pub fn calling_convention(mut self, cc: AmdgpuCallingConvention) -> Self {
        self.calling_convention = cc;
        self
    }

    pub fn simt_emulation(mut self, enable: bool) -> Self {
        self.emit_simt = enable;
        self
    }

    pub fn scalarize_uniform(mut self, enable: bool) -> Self {
        self.scalarize = enable;
        self
    }

    pub fn add_kernel_arg(mut self, arg: AmdgpuKernelArg) -> Self {
        self.kernel_args.push(arg);
        self
    }

    pub fn build(self) -> AMDGPUX86Bridge {
        let mut bridge = AMDGPUX86Bridge::new(self.isa);
        bridge.set_opt_level(self.opt_level);
        bridge.set_calling_convention(self.calling_convention);

        if self.use_avx512 {
            bridge.enable_avx512();
        } else {
            bridge.target_machine.use_avx = self.use_avx;
            bridge.target_machine.use_avx2 = self.use_avx2;
        }
        bridge.target_machine.emit_simt_masks = self.emit_simt;
        bridge.cross_target.scalarize_uniform = self.scalarize;
        bridge.cross_target.emit_lane_masks = self.emit_simt;

        for arg in self.kernel_args {
            bridge.add_kernel_arg(arg);
        }

        bridge
    }
}

// ═══════════════════════════════════════════════════════════════════════════════
// Test Suite
// ═══════════════════════════════════════════════════════════════════════════════

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

    // ── X86 Register Tests ─────────────────────────────────────────────────

    #[test]
    fn test_x86_reg_names() {
        assert_eq!(X86PhysReg::RAX.name(), "rax");
        assert_eq!(X86PhysReg::RBX.name(), "rbx");
        assert_eq!(X86PhysReg::XMM0.name(), "xmm0");
        assert_eq!(X86PhysReg::YMM0.name(), "ymm0");
        assert_eq!(X86PhysReg::ZMM0.name(), "zmm0");
    }

    #[test]
    fn test_x86_reg_sizes() {
        assert_eq!(X86PhysReg::RAX.size_bits(), 64);
        assert_eq!(X86PhysReg::EAX.size_bits(), 32);
        assert_eq!(X86PhysReg::AX.size_bits(), 16);
        assert_eq!(X86PhysReg::AL.size_bits(), 8);
        assert_eq!(X86PhysReg::XMM0.size_bits(), 128);
        assert_eq!(X86PhysReg::YMM0.size_bits(), 256);
        assert_eq!(X86PhysReg::ZMM0.size_bits(), 512);
    }

    #[test]
    fn test_x86_reg_classes() {
        assert_eq!(X86PhysReg::RAX.reg_class(), X86RegClass::GPR);
        assert_eq!(X86PhysReg::XMM0.reg_class(), X86RegClass::XMM);
        assert_eq!(X86PhysReg::YMM0.reg_class(), X86RegClass::YMM);
        assert_eq!(X86PhysReg::ZMM0.reg_class(), X86RegClass::ZMM);
        assert_eq!(X86PhysReg::K0.reg_class(), X86RegClass::Mask);
        assert_eq!(X86PhysReg::EFLAGS.reg_class(), X86RegClass::Flags);
    }

    #[test]
    fn test_x86_reg_encoding() {
        assert_eq!(X86PhysReg::RAX.encoding(), 0);
        assert_eq!(X86PhysReg::RCX.encoding(), 1);
        assert_eq!(X86PhysReg::RDX.encoding(), 2);
        assert_eq!(X86PhysReg::RBX.encoding(), 3);
    }

    #[test]
    fn test_x86_reg_display() {
        assert_eq!(format!("{}", X86PhysReg::RAX), "rax");
        assert_eq!(format!("{}", X86PhysReg::XMM15), "xmm15");
    }

    // ── X86 Opcode Tests ───────────────────────────────────────────────────

    #[test]
    fn test_x86_opcode_mnemonics() {
        assert_eq!(X86Opcode::Add.mnemonic(), "add");
        assert_eq!(X86Opcode::Sub.mnemonic(), "sub");
        assert_eq!(X86Opcode::Mul.mnemonic(), "mul");
        assert_eq!(X86Opcode::And.mnemonic(), "and");
        assert_eq!(X86Opcode::Or.mnemonic(), "or");
        assert_eq!(X86Opcode::Xor.mnemonic(), "xor");
        assert_eq!(X86Opcode::Nop.mnemonic(), "nop");
        assert_eq!(X86Opcode::Jmp.mnemonic(), "jmp");
        assert_eq!(X86Opcode::Ret.mnemonic(), "ret");
    }

    #[test]
    fn test_x86_sse_opcode_mnemonics() {
        assert_eq!(X86Opcode::Addss.mnemonic(), "addss");
        assert_eq!(X86Opcode::Subss.mnemonic(), "subss");
        assert_eq!(X86Opcode::Mulss.mnemonic(), "mulss");
        assert_eq!(X86Opcode::Divss.mnemonic(), "divss");
        assert_eq!(X86Opcode::Maxss.mnemonic(), "maxss");
        assert_eq!(X86Opcode::Minss.mnemonic(), "minss");
    }

    // ── Register Mapping Tests ─────────────────────────────────────────────

    #[test]
    fn test_register_mapping_new() {
        let mapping = RegisterMapping::new(0, "sgpr0", AmdgpuRegClass::SGPR32, X86PhysReg::RAX);
        assert_eq!(mapping.amdgpu_reg_id, 0);
        assert_eq!(mapping.amdgpu_reg_name, "sgpr0");
        assert_eq!(mapping.x86_reg, X86PhysReg::RAX);
        assert!(mapping.is_direct);
        assert!(!mapping.requires_swizzle);
    }

    // ── Instruction Mapping Tests ──────────────────────────────────────────

    #[test]
    fn test_instruction_mapping_direct() {
        let mapping = InstructionMapping::new(AmdgpuOpcode::SAddU32, vec![X86Opcode::Add], 1);
        assert!(mapping.is_direct);
        assert_eq!(mapping.expansion_factor, 1);
        assert!(mapping.preserves_semantics);
    }

    #[test]
    fn test_instruction_mapping_expanded() {
        let mapping = InstructionMapping::new(
            AmdgpuOpcode::SAbsI32,
            vec![X86Opcode::MovRR, X86Opcode::Neg, X86Opcode::CmovL],
            3,
        );
        assert!(!mapping.is_direct);
        assert_eq!(mapping.expansion_factor, 3);
    }

    #[test]
    fn test_instruction_mapping_with_simt() {
        let mapping = InstructionMapping::new(AmdgpuOpcode::SCmpEqI32, vec![X86Opcode::Cmp], 1)
            .with_simt(true);
        assert!(mapping.simt_aware);
    }

    // ── Calling Convention Tests ───────────────────────────────────────────

    #[test]
    fn test_kernel_calling_convention() {
        let cc = AmdgpuCallingConvention::Kernel;
        assert!(cc.uses_sgpr_args());
        assert!(!cc.returns_via_vgpr());
        assert!(!cc.supports_sret());
        assert_eq!(cc.implicit_sgpr_count(), 6);
        assert_eq!(cc.implicit_vgpr_count(), 1);
    }

    #[test]
    fn test_device_calling_convention() {
        let cc = AmdgpuCallingConvention::DeviceFunction;
        assert!(!cc.uses_sgpr_args());
        assert!(cc.returns_via_vgpr());
        assert!(cc.supports_sret());
        assert_eq!(cc.implicit_sgpr_count(), 0);
        assert_eq!(cc.implicit_vgpr_count(), 0);
    }

    #[test]
    fn test_fast_calling_convention() {
        let cc = AmdgpuCallingConvention::Fast;
        assert!(cc.uses_sgpr_args());
        assert!(cc.returns_via_vgpr());
        assert!(!cc.supports_sret());
    }

    #[test]
    fn test_indirect_calling_convention() {
        let cc = AmdgpuCallingConvention::Indirect;
        assert!(!cc.uses_sgpr_args());
        assert!(cc.returns_via_vgpr());
        assert!(cc.supports_sret());
    }

    #[test]
    fn test_calling_convention_display() {
        assert_eq!(AmdgpuCallingConvention::Kernel.to_string(), "amdgpu_kernel");
        assert_eq!(
            AmdgpuCallingConvention::DeviceFunction.to_string(),
            "amdgpu_device"
        );
    }

    // ── Address Space Tests ────────────────────────────────────────────────

    #[test]
    fn test_addr_space_from_u32() {
        assert_eq!(AmdgpuAddrSpace::from_u32(1), Some(AmdgpuAddrSpace::Global));
        assert_eq!(AmdgpuAddrSpace::from_u32(3), Some(AmdgpuAddrSpace::Local));
        assert_eq!(
            AmdgpuAddrSpace::from_u32(4),
            Some(AmdgpuAddrSpace::Constant)
        );
        assert_eq!(AmdgpuAddrSpace::from_u32(5), Some(AmdgpuAddrSpace::Private));
        assert_eq!(AmdgpuAddrSpace::from_u32(0), Some(AmdgpuAddrSpace::Generic));
        assert_eq!(AmdgpuAddrSpace::from_u32(99), None);
    }

    #[test]
    fn test_addr_space_names() {
        assert_eq!(AmdgpuAddrSpace::Global.name(), "global");
        assert_eq!(AmdgpuAddrSpace::Local.name(), "local");
        assert_eq!(AmdgpuAddrSpace::Private.name(), "private");
        assert_eq!(AmdgpuAddrSpace::Constant.name(), "constant");
    }

    #[test]
    fn test_addr_space_to_llvm() {
        assert_eq!(AmdgpuAddrSpace::Global.to_llvm_addrspace(), 1);
        assert_eq!(AmdgpuAddrSpace::Local.to_llvm_addrspace(), 3);
    }

    // ── Kernel Arg Tests ───────────────────────────────────────────────────

    #[test]
    fn test_kernel_arg_creation() {
        let arg = AmdgpuKernelArg::new("input", 4, 4, 0);
        assert_eq!(arg.name, "input");
        assert_eq!(arg.size_bytes, 4);
        assert_eq!(arg.alignment, 4);
        assert_eq!(arg.offset, 0);
        assert!(!arg.is_pointer);
        assert_eq!(arg.addr_space, AmdgpuAddrSpace::Global);
    }

    #[test]
    fn test_kernel_arg_with_pointer() {
        let arg = AmdgpuKernelArg::new("ptr", 8, 8, 0).with_pointer(true);
        assert!(arg.is_pointer);
        assert_eq!(arg.size_bytes, 8);
    }

    #[test]
    fn test_kernel_arg_with_addr_space() {
        let arg = AmdgpuKernelArg::new("lds_data", 4, 4, 0).with_addr_space(AmdgpuAddrSpace::Local);
        assert_eq!(arg.addr_space, AmdgpuAddrSpace::Local);
    }

    #[test]
    fn test_kernel_arg_with_registers() {
        let arg = AmdgpuKernelArg::new("val", 4, 4, 0)
            .with_sgpr(0)
            .with_vgpr(1);
        assert_eq!(arg.sgpr_index, Some(0));
        assert_eq!(arg.vgpr_index, Some(1));
    }

    // ── Register Info X86 Tests ────────────────────────────────────────────

    #[test]
    fn test_register_info_x86_creation() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX900);
        assert!(info.mapped_count() > 0);
        assert_eq!(info.available_vgprs, 256);
        assert!(!info.has_agprs); // GFX900 does not have AGPRs
    }

    #[test]
    fn test_register_info_x86_rdna() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX1010);
        assert!(info.mapped_count() > 0);
        // AGPRs available on RDNA
    }

    #[test]
    fn test_register_info_x86_agpr() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX908);
        assert!(info.has_agprs);
    }

    #[test]
    fn test_register_info_map_to_x86_sgpr() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX900);
        let x86_reg = info.map_to_x86(0); // SGPR0 -> RAX
        assert_eq!(x86_reg, Some(X86PhysReg::RAX));
    }

    #[test]
    fn test_register_info_map_to_x86_vgpr() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX900);
        let x86_reg = info.map_to_x86(256); // VGPR0 -> XMM0
        assert_eq!(x86_reg, Some(X86PhysReg::XMM0));
    }

    #[test]
    fn test_register_info_map_from_x86() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX900);
        let amdgpu_id = info.map_from_x86(X86PhysReg::RAX);
        assert_eq!(amdgpu_id, Some(0));
    }

    #[test]
    fn test_register_info_map_from_xmm() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX900);
        let amdgpu_id = info.map_from_x86(X86PhysReg::XMM0);
        assert_eq!(amdgpu_id, Some(256));
    }

    #[test]
    fn test_register_info_map_nonexistent() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX900);
        // Unmapped register
        assert_eq!(info.map_to_x86(9999), None);
        assert_eq!(info.map_from_x86(X86PhysReg::ST0), None);
    }

    // ── Target Machine Tests ───────────────────────────────────────────────

    #[test]
    fn test_amdgpu_x86_target_machine_creation() {
        let tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        assert!(tm.base.is_gcn());
        assert_eq!(tm.x86_triple, "x86_64-unknown-linux-gnu");
        assert!(tm.mapping_count() > 0);
    }

    #[test]
    fn test_amdgpu_x86_target_machine_from_cpu() {
        let tm = AmdgpuX86TargetMachine::from_cpu("gfx900").unwrap();
        assert!(tm.base.is_gcn());
    }

    #[test]
    fn test_amdgpu_x86_target_machine_rdna() {
        let tm = AmdgpuX86TargetMachine::from_cpu("gfx1010").unwrap();
        assert!(tm.base.is_rdna());
    }

    #[test]
    fn test_amdgpu_x86_target_machine_invalid_cpu() {
        let tm = AmdgpuX86TargetMachine::from_cpu("invalid");
        assert!(tm.is_none());
    }

    #[test]
    fn test_target_machine_lookup_mapping() {
        let tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        let mapping = tm.lookup_mapping(AmdgpuOpcode::SAddU32);
        assert!(mapping.is_some());
        let mapping = mapping.unwrap();
        assert_eq!(mapping.x86_opcodes.len(), 1);
        assert_eq!(mapping.x86_opcodes[0], X86Opcode::Add);
    }

    #[test]
    fn test_target_machine_lookup_nonexistent() {
        let tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        let mapping = tm.lookup_mapping(AmdgpuOpcode::SWaitcnt); // Not mapped
        assert!(mapping.is_none());
    }

    #[test]
    fn test_target_machine_all_mappings() {
        let tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        let all = tm.all_mappings();
        assert!(all.len() >= 80); // Should have 80+ mappings
    }

    #[test]
    fn test_target_machine_vgpr_to_avx() {
        let mut tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        // Default: vgpr0 → xmm0
        let reg = tm.map_vgpr_to_x86_vector(256);
        assert_eq!(reg, Some(X86PhysReg::XMM0));
    }

    #[test]
    fn test_target_machine_vgpr_to_avx512() {
        let mut tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        tm.enable_avx512();
        let reg = tm.map_vgpr_to_x86_vector(256);
        assert_eq!(reg, Some(X86PhysReg::ZMM0));
    }

    #[test]
    fn test_target_machine_opt_level() {
        let mut tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        assert_eq!(tm.opt_level, 2);
        tm.set_opt_level(3);
        assert_eq!(tm.opt_level, 3);
        tm.set_opt_level(5);
        assert_eq!(tm.opt_level, 3); // Clamped to 3
    }

    #[test]
    fn test_target_machine_x86_features() {
        let tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        let features = tm.x86_feature_string();
        assert!(features.contains("sse2"));
        assert!(features.contains("avx"));
        assert!(features.contains("avx2"));
        assert!(!features.contains("avx512f"));
    }

    #[test]
    fn test_target_machine_x86_features_avx512() {
        let mut tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        tm.enable_avx512();
        let features = tm.x86_feature_string();
        assert!(features.contains("avx512f"));
        assert!(features.contains("avx512dq"));
        assert!(features.contains("avx512vl"));
    }

    #[test]
    fn test_target_machine_data_layout() {
        let tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        let dl = tm.get_x86_data_layout();
        assert!(dl.contains("e-m:e"));
        assert!(dl.starts_with("e-m:e"));
    }

    // ── Bridge Tests ───────────────────────────────────────────────────────

    #[test]
    fn test_bridge_creation() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        assert!(bridge.is_ready());
        assert_eq!(bridge.calling_convention, AmdgpuCallingConvention::Kernel);
        assert!(bridge.kernel_args.is_empty());
    }

    #[test]
    fn test_bridge_from_cpu() {
        let bridge = AMDGPUX86Bridge::from_cpu("gfx1030").unwrap();
        assert!(bridge.supports_wave32());
        assert_eq!(bridge.isa_version(), AmdgpuIsaVersion::GFX1030);
    }

    #[test]
    fn test_bridge_from_cpu_invalid() {
        let bridge = AMDGPUX86Bridge::from_cpu("invalid");
        assert!(bridge.is_none());
    }

    #[test]
    fn test_bridge_calling_convention() {
        let mut bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        bridge.set_calling_convention(AmdgpuCallingConvention::DeviceFunction);
        assert_eq!(
            bridge.calling_convention,
            AmdgpuCallingConvention::DeviceFunction
        );
    }

    #[test]
    fn test_bridge_kernel_args() {
        let mut bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let arg = AmdgpuKernelArg::new("a", 4, 4, 0);
        bridge.add_kernel_arg(arg);
        assert_eq!(bridge.kernel_args.len(), 1);

        let arg2 = AmdgpuKernelArg::new("b", 8, 8, 4);
        bridge.add_kernel_arg(arg2);
        assert_eq!(bridge.kernel_args.len(), 2);

        bridge.clear_kernel_args();
        assert_eq!(bridge.kernel_args.len(), 0);
    }

    #[test]
    fn test_bridge_translate_instruction() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let result = bridge.translate_instruction(AmdgpuOpcode::SAddU32);
        assert!(result.is_some());
        let (x86_ops, expansion) = result.unwrap();
        assert_eq!(expansion, 1);
        assert_eq!(x86_ops.len(), 1);
        assert_eq!(x86_ops[0], X86Opcode::Add);
    }

    #[test]
    fn test_bridge_translate_expanded() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let result = bridge.translate_instruction(AmdgpuOpcode::SMovB64);
        assert!(result.is_some());
        let (_x86_ops, expansion) = result.unwrap();
        assert_eq!(expansion, 2);
    }

    #[test]
    fn test_bridge_direct_mapping() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        assert!(bridge.has_direct_mapping(AmdgpuOpcode::SAddU32));
        assert!(!bridge.has_direct_mapping(AmdgpuOpcode::SMovB64));
    }

    #[test]
    fn test_bridge_compute_expansion() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let ops = &[
            AmdgpuOpcode::SAddU32, // 1
            AmdgpuOpcode::SMovB64, // 2
            AmdgpuOpcode::SMulI32, // 1
        ];
        let expansion = bridge.compute_expansion(ops);
        assert_eq!(expansion, 4);
    }

    #[test]
    fn test_bridge_avx512() {
        let mut bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        bridge.enable_avx512();
        assert!(bridge.stats.avx512_enabled);
        let features = bridge.x86_feature_string();
        assert!(features.contains("avx512f"));
    }

    #[test]
    fn test_bridge_opt_level() {
        let mut bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        bridge.set_opt_level(3);
        assert_eq!(bridge.target_machine.opt_level, 3);
    }

    #[test]
    fn test_bridge_x86_prologue() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let prologue = bridge.emit_x86_prologue();
        assert!(prologue.contains("AMDGPU Kernel Prologue"));
        assert!(prologue.contains("gfx900"));
        assert!(prologue.contains("amdgpu_kernel"));
    }

    #[test]
    fn test_bridge_x86_epilogue() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let epilogue = bridge.emit_x86_epilogue();
        assert!(epilogue.contains("AMDGPU Kernel Epilogue"));
    }

    #[test]
    fn test_bridge_x86_prologue_with_args() {
        let mut bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        bridge.add_kernel_arg(AmdgpuKernelArg::new("input", 4, 4, 0));
        let prologue = bridge.emit_x86_prologue();
        assert!(prologue.contains("input"));
        assert!(prologue.contains("Kernel Launch Parameters"));
    }

    #[test]
    fn test_bridge_get_triple() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let triple = bridge.get_amdgpu_triple();
        assert!(triple.contains("amdgcn"));
    }

    #[test]
    fn test_bridge_x86_data_layout() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let dl = bridge.get_x86_data_layout();
        assert!(!dl.is_empty());
    }

    #[test]
    fn test_bridge_suggest_register_allocation() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let regs = bridge.suggest_x86_register_allocation(8);
        assert_eq!(regs.len(), 8);
    }

    #[test]
    fn test_bridge_suggest_register_allocation_max() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let regs = bridge.suggest_x86_register_allocation(32);
        assert_eq!(regs.len(), 16); // Capped at 16
    }

    #[test]
    fn test_bridge_divergence_analysis() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let ops = &[
            AmdgpuOpcode::SAddU32, // scalar
            AmdgpuOpcode::VAddF32, // vector
            AmdgpuOpcode::VMulF32, // vector
            AmdgpuOpcode::SAndB32, // scalar
            AmdgpuOpcode::SOrB32,  // scalar
            AmdgpuOpcode::VSubF32, // vector
        ];
        let analysis = bridge.analyze_divergence(ops);
        assert_eq!(analysis.scalar_count, 3);
        assert_eq!(analysis.vector_count, 3);
    }

    #[test]
    fn test_bridge_analyze_divergence_all_scalar() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let ops = &[
            AmdgpuOpcode::SAddU32,
            AmdgpuOpcode::SAndB32,
            AmdgpuOpcode::SOrB32,
        ];
        let analysis = bridge.analyze_divergence(ops);
        assert_eq!(analysis.scalar_count, 3);
        assert_eq!(analysis.vector_count, 0);
    }

    #[test]
    fn test_bridge_analyze_divergence_all_vector() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let ops = &[AmdgpuOpcode::VAddF32, AmdgpuOpcode::VMulF32];
        let analysis = bridge.analyze_divergence(ops);
        assert_eq!(analysis.scalar_count, 0);
        assert_eq!(analysis.vector_count, 2);
    }

    #[test]
    fn test_bridge_stats() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let stats = bridge.get_stats();
        assert_eq!(stats.instructions_translated, 0);
        assert_eq!(stats.kernels_processed, 0);
        assert!(!stats.avx512_enabled);
    }

    #[test]
    fn test_bridge_stats_after_avx512() {
        let mut bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        bridge.enable_avx512();
        let stats = bridge.get_stats();
        assert!(stats.avx512_enabled);
    }

    #[test]
    fn test_bridge_reset_stats() {
        let mut bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        bridge.enable_avx512();
        bridge.reset_stats();
        let stats = bridge.get_stats();
        assert!(!stats.avx512_enabled);
    }

    // ── Cross-Target Tests ─────────────────────────────────────────────────

    #[test]
    fn test_cross_target_creation() {
        let ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        assert_eq!(ct.wavefront_size, 64);
        assert_eq!(ct.max_vgprs, 256);
        assert_eq!(ct.max_sgprs, 104);
        assert!(ct.scalarize_uniform);
        assert!(ct.emit_lane_masks);
    }

    #[test]
    fn test_cross_target_wave32() {
        let ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX1010);
        assert_eq!(ct.wavefront_size, 32);
    }

    #[test]
    fn test_cross_target_allocate_vgpr() {
        let mut ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        let reg = ct.allocate_vgpr().unwrap();
        assert!(reg < 256);
        // Second allocation should get next available
        let reg2 = ct.allocate_vgpr().unwrap();
        assert!(reg2 != reg);
    }

    #[test]
    fn test_cross_target_allocate_sgpr() {
        let mut ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        let reg = ct.allocate_sgpr().unwrap();
        assert!(reg < 104);
    }

    #[test]
    fn test_cross_target_free_vgpr() {
        let mut ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        let reg = ct.allocate_vgpr().unwrap();
        assert!(ct.allocated_vgprs.contains(&reg));
        ct.free_vgpr(reg);
        assert!(!ct.allocated_vgprs.contains(&reg));
    }

    #[test]
    fn test_cross_target_free_sgpr() {
        let mut ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        let reg = ct.allocate_sgpr().unwrap();
        ct.free_sgpr(reg);
        assert!(!ct.allocated_sgprs.contains(&reg));
    }

    #[test]
    fn test_cross_target_lane_reg() {
        let mut ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        ct.set_lane_id_reg(X86PhysReg::RAX);
        assert_eq!(ct.lane_id_reg, Some(X86PhysReg::RAX));
        ct.set_exec_mask_reg(X86PhysReg::RBX);
        assert_eq!(ct.exec_mask_reg, Some(X86PhysReg::RBX));
    }

    #[test]
    fn test_is_scalar_format() {
        assert!(CrossTargetAmdgpu::is_scalar_format(
            AmdgpuEncodingFormat::SOP1
        ));
        assert!(CrossTargetAmdgpu::is_scalar_format(
            AmdgpuEncodingFormat::SOP2
        ));
        assert!(CrossTargetAmdgpu::is_scalar_format(
            AmdgpuEncodingFormat::SOPP
        ));
        assert!(CrossTargetAmdgpu::is_scalar_format(
            AmdgpuEncodingFormat::SMEM
        ));
        assert!(!CrossTargetAmdgpu::is_scalar_format(
            AmdgpuEncodingFormat::VOP1
        ));
    }

    #[test]
    fn test_is_vector_format() {
        assert!(CrossTargetAmdgpu::is_vector_format(
            AmdgpuEncodingFormat::VOP1
        ));
        assert!(CrossTargetAmdgpu::is_vector_format(
            AmdgpuEncodingFormat::VOP2
        ));
        assert!(CrossTargetAmdgpu::is_vector_format(
            AmdgpuEncodingFormat::FLAT
        ));
        assert!(CrossTargetAmdgpu::is_vector_format(
            AmdgpuEncodingFormat::DS
        ));
        assert!(CrossTargetAmdgpu::is_vector_format(
            AmdgpuEncodingFormat::MIMG
        ));
        assert!(!CrossTargetAmdgpu::is_vector_format(
            AmdgpuEncodingFormat::SOP1
        ));
    }

    #[test]
    fn test_cross_target_emit_lane_init() {
        let ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        let code = ct.emit_lane_init();
        assert!(code.contains("SIMT Lane Mask Initialization"));
        assert!(code.contains("64"));
    }

    #[test]
    fn test_cross_target_emit_exec_predicate() {
        let ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        let code = ct.emit_exec_predicate(false);
        assert!(code.contains("exec_mask"));
    }

    #[test]
    fn test_cross_target_reset() {
        let mut ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        ct.allocate_vgpr().unwrap();
        ct.allocate_sgpr().unwrap();
        ct.set_lane_id_reg(X86PhysReg::RAX);
        assert_eq!(ct.allocated_count(), 2);
        ct.reset_allocations();
        assert_eq!(ct.allocated_count(), 0);
        assert_eq!(ct.lane_id_reg, None);
    }

    #[test]
    fn test_cross_target_lanes_per_wavefront() {
        let ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        assert_eq!(ct.lanes_per_wavefront(), 64);
    }

    // ── Kernel Descriptor Tests ────────────────────────────────────────────

    #[test]
    fn test_kernel_descriptor_new() {
        let kd = AmdgpuKernelDescriptorX86::new("test_kernel");
        assert_eq!(kd.name, "test_kernel");
        assert_eq!(kd.sgpr_count, 0);
        assert_eq!(kd.vgpr_count, 0);
        assert_eq!(kd.group_segment_size, 0);
        assert!(!kd.uses_dynamic_lds);
        assert_eq!(kd.kernel_code_entry_byte_offset, 256);
    }

    #[test]
    fn test_kernel_descriptor_with_args() {
        let mut kd = AmdgpuKernelDescriptorX86::new("kernel_with_args");
        kd.add_arg(AmdgpuKernelArg::new("a", 4, 4, 0));
        kd.add_arg(AmdgpuKernelArg::new("b", 8, 8, 4));
        assert_eq!(kd.args.len(), 2);
    }

    #[test]
    fn test_kernel_descriptor_register_usage() {
        let mut kd = AmdgpuKernelDescriptorX86::new("kernel");
        kd.set_register_usage(16, 32);
        assert_eq!(kd.sgpr_count, 16);
        assert_eq!(kd.vgpr_count, 32);
    }

    #[test]
    fn test_kernel_descriptor_spill() {
        let mut kd = AmdgpuKernelDescriptorX86::new("kernel");
        kd.set_spill(64, 128);
        assert_eq!(kd.sgpr_spill_size, 64);
        assert_eq!(kd.vgpr_spill_size, 128);
    }

    #[test]
    fn test_kernel_descriptor_emit_launch_stub() {
        let mut kd = AmdgpuKernelDescriptorX86::new("my_kernel");
        kd.set_register_usage(10, 20);
        let stub = kd.emit_x86_launch_stub();
        assert!(stub.contains("my_kernel"));
        assert!(stub.contains("X86 Launch Stub"));
    }

    #[test]
    fn test_kernel_descriptor_metadata() {
        let mut kd = AmdgpuKernelDescriptorX86::new("test");
        kd.set_register_usage(8, 16);
        kd.set_spill(0, 0);
        kd.group_segment_size = 4096;
        let md = kd.to_metadata_string();
        assert!(md.contains("\"name\":\"test\""));
        assert!(md.contains("\"sgpr_count\":8"));
        assert!(md.contains("\"vgpr_count\":16"));
        assert!(md.contains("\"group_segment_size\":4096"));
    }

    // ── ABI Context Tests ──────────────────────────────────────────────────

    #[test]
    fn test_abi_context_new() {
        let ctx = AmdgpuKernelAbiContext::new("kernel_abi");
        assert_eq!(ctx.kernel_name, "kernel_abi");
        assert_eq!(ctx.calling_convention, AmdgpuCallingConvention::Kernel);
        assert!(ctx.local_id_x.is_none());
    }

    #[test]
    fn test_abi_context_emit_workitem_id() {
        let ctx = AmdgpuKernelAbiContext::new("kernel");
        let code = ctx.emit_load_workitem_id_x86();
        assert!(code.contains("Work-Item ID"));
        assert!(code.contains("VGPR0"));
    }

    #[test]
    fn test_abi_context_emit_global_id() {
        let ctx = AmdgpuKernelAbiContext::new("kernel");
        let code = ctx.emit_compute_global_id_x();
        assert!(code.contains("global_id.x"));
    }

    #[test]
    fn test_abi_context_emit_aql_dispatch() {
        let ctx = AmdgpuKernelAbiContext::new("kernel");
        let code = ctx.emit_aql_dispatch();
        assert!(code.contains("AQL Dispatch Packet"));
        assert!(code.contains("grid_size_x"));
    }

    // ── Memory Model Tests ─────────────────────────────────────────────────

    #[test]
    fn test_memory_order_relaxed() {
        let order = AmdgpuMemoryOrder::Relaxed;
        assert!(!order.to_x86_mfence());
        assert!(!order.to_x86_lfence());
        assert!(!order.to_x86_sfence());
    }

    #[test]
    fn test_memory_order_seq_cst() {
        let order = AmdgpuMemoryOrder::SequentiallyConsistent;
        assert!(order.to_x86_mfence());
    }

    #[test]
    fn test_memory_order_acquire() {
        let order = AmdgpuMemoryOrder::Acquire;
        assert!(order.to_x86_lfence());
        assert!(!order.to_x86_sfence());
    }

    #[test]
    fn test_memory_order_release() {
        let order = AmdgpuMemoryOrder::Release;
        assert!(!order.to_x86_lfence());
        assert!(order.to_x86_sfence());
    }

    #[test]
    fn test_memory_order_acq_rel() {
        let order = AmdgpuMemoryOrder::AcquireRelease;
        assert!(order.to_x86_lfence());
        assert!(order.to_x86_sfence());
    }

    #[test]
    fn test_memory_scope_workitem() {
        let scope = AmdgpuMemoryScope::WorkItem;
        assert!(!scope.requires_barrier());
        assert_eq!(scope.to_x86_fence(), None);
    }

    #[test]
    fn test_memory_scope_workgroup() {
        let scope = AmdgpuMemoryScope::Workgroup;
        assert!(scope.requires_barrier());
        assert_eq!(scope.to_x86_fence(), Some("__syncthreads()"));
    }

    #[test]
    fn test_memory_scope_device() {
        let scope = AmdgpuMemoryScope::Device;
        assert!(scope.requires_barrier());
        assert_eq!(scope.to_x86_fence(), Some("mfence"));
    }

    #[test]
    fn test_memory_scope_system() {
        let scope = AmdgpuMemoryScope::System;
        assert!(scope.requires_barrier());
        assert_eq!(scope.to_x86_fence(), Some("mfence"));
    }

    // ── Wavefront Emulator Tests ───────────────────────────────────────────

    #[test]
    fn test_wavefront_emulator_new() {
        let we = WavefrontEmulator::new(64);
        assert_eq!(we.wave_size, 64);
        assert_eq!(we.num_waves, 1);
        assert!(we.all_lanes_active());
    }

    #[test]
    fn test_wavefront_emulator_wave32() {
        let we = WavefrontEmulator::new(32);
        assert_eq!(we.wave_size, 32);
        assert_eq!(we.active_lane_count(), 32);
    }

    #[test]
    fn test_wavefront_emulator_exec_mask() {
        let mut we = WavefrontEmulator::new(64);
        we.set_exec_mask(0x0000_0000_0000_00FF);
        assert_eq!(we.active_lane_count(), 8);
        assert!(!we.all_lanes_active());
        assert!(we.lane_active(0));
        assert!(we.lane_active(7));
        assert!(!we.lane_active(8));
    }

    #[test]
    fn test_wavefront_emulator_lane_active() {
        let mut we = WavefrontEmulator::new(64);
        we.set_exec_mask(1 << 5);
        assert!(!we.lane_active(0));
        assert!(we.lane_active(5));
        assert!(!we.lane_active(6));
    }

    #[test]
    fn test_wavefront_emulator_loop_header() {
        let we = WavefrontEmulator::new(64);
        let code = we.emit_simt_loop_header();
        assert!(code.contains("SIMT Wavefront Loop"));
        assert!(code.contains("64 lanes"));
    }

    #[test]
    fn test_wavefront_emulator_loop_footer() {
        let we = WavefrontEmulator::new(64);
        let code = we.emit_simt_loop_footer();
        assert!(code.contains(".L_lane_inactive"));
        assert!(code.contains(".L_simt_done"));
    }

    #[test]
    fn test_lane_mask() {
        assert_eq!(WavefrontEmulator::lane_mask(0, 4), 0xF);
        assert_eq!(WavefrontEmulator::lane_mask(4, 4), 0xF0);
        assert_eq!(WavefrontEmulator::lane_mask(0, 0), 0);
        assert_eq!(WavefrontEmulator::lane_mask(60, 4), 0xF000_0000_0000_0000);
    }

    // ── Bridge Builder Tests ───────────────────────────────────────────────

    #[test]
    fn test_builder_default() {
        let builder = AmdgpuBridgeBuilder::new(AmdgpuIsaVersion::GFX900);
        let bridge = builder.build();
        assert!(bridge.is_ready());
        assert_eq!(bridge.target_machine.opt_level, 2);
        assert!(bridge.target_machine.use_avx);
    }

    #[test]
    fn test_builder_with_options() {
        let builder = AmdgpuBridgeBuilder::new(AmdgpuIsaVersion::GFX1010)
            .opt_level(3)
            .avx512(true)
            .simt_emulation(false)
            .calling_convention(AmdgpuCallingConvention::DeviceFunction);
        let bridge = builder.build();
        assert_eq!(bridge.target_machine.opt_level, 3);
        assert!(bridge.target_machine.use_avx512);
        assert!(!bridge.target_machine.emit_simt_masks);
        assert_eq!(
            bridge.calling_convention,
            AmdgpuCallingConvention::DeviceFunction
        );
    }

    #[test]
    fn test_builder_with_kernel_args() {
        let builder = AmdgpuBridgeBuilder::new(AmdgpuIsaVersion::GFX900)
            .add_kernel_arg(AmdgpuKernelArg::new("a", 4, 4, 0))
            .add_kernel_arg(AmdgpuKernelArg::new("b", 8, 8, 4));
        let bridge = builder.build();
        assert_eq!(bridge.kernel_args.len(), 2);
    }

    #[test]
    fn test_builder_from_cpu() {
        let builder = AmdgpuBridgeBuilder::from_cpu("gfx1030").unwrap();
        let bridge = builder.build();
        assert!(bridge.supports_wave32());
    }

    #[test]
    fn test_builder_from_cpu_invalid() {
        let builder = AmdgpuBridgeBuilder::from_cpu("invalid");
        assert!(builder.is_none());
    }

    #[test]
    fn test_builder_avx() {
        let mut bridge = AmdgpuBridgeBuilder::new(AmdgpuIsaVersion::GFX900)
            .avx(false)
            .avx2(false)
            .build();
        // Should still build, XMM is the fallback
        let regs = bridge.suggest_x86_register_allocation(1);
        assert_eq!(regs.len(), 1);
    }

    // ── Divergence Analysis Tests ──────────────────────────────────────────

    #[test]
    fn test_divergence_analysis_default() {
        let da = DivergenceAnalysis::default();
        assert_eq!(da.scalar_count, 0);
        assert_eq!(da.vector_count, 0);
        assert_eq!(da.total(), 0);
        assert_eq!(da.scalar_ratio(), 0.0);
        assert_eq!(da.vector_ratio(), 0.0);
    }

    #[test]
    fn test_divergence_analysis_ratios() {
        let da = DivergenceAnalysis {
            scalar_count: 30,
            vector_count: 70,
            total_count: 100,
        };
        assert!((da.scalar_ratio() - 0.3).abs() < 0.001);
        assert!((da.vector_ratio() - 0.7).abs() < 0.001);
    }

    // ── Edge Cases and Coverage ────────────────────────────────────────────

    #[test]
    fn test_empty_bridge_translate() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        // Unmapped opcode should return None
        let result = bridge.translate_instruction(AmdgpuOpcode::SWaitcnt);
        assert!(result.is_none());
    }

    #[test]
    fn test_empty_expansion() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let expansion = bridge.compute_expansion(&[]);
        assert_eq!(expansion, 0);
    }

    #[test]
    fn test_all_isa_versions_build() {
        for isa in &[
            AmdgpuIsaVersion::GFX600,
            AmdgpuIsaVersion::GFX700,
            AmdgpuIsaVersion::GFX800,
            AmdgpuIsaVersion::GFX900,
            AmdgpuIsaVersion::GFX906,
            AmdgpuIsaVersion::GFX908,
            AmdgpuIsaVersion::GFX90A,
            AmdgpuIsaVersion::GFX1010,
            AmdgpuIsaVersion::GFX1030,
            AmdgpuIsaVersion::GFX1100,
            AmdgpuIsaVersion::GFX1150,
        ] {
            let bridge = AMDGPUX86Bridge::new(*isa);
            assert!(bridge.is_ready());
        }
    }

    #[test]
    fn test_instruction_mapping_count_at_least_80() {
        let tm = AmdgpuX86TargetMachine::new(AmdgpuIsaVersion::GFX900);
        assert!(tm.mapping_count() >= 80);
    }

    #[test]
    fn test_suggest_x86_regs_zero_count() {
        let bridge = AMDGPUX86Bridge::new(AmdgpuIsaVersion::GFX900);
        let regs = bridge.suggest_x86_register_allocation(0);
        assert!(regs.is_empty());
    }

    #[test]
    fn test_register_info_with_agpr() {
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX90A);
        assert!(info.has_agprs);
    }

    #[test]
    fn test_nv_reg_mapping_complete() {
        // Verify all 16 SGPRs + 16 VGPRs have mappings
        let info = AmdgpuRegisterInfoX86::new(AmdgpuIsaVersion::GFX1010);
        assert_eq!(info.mapped_count(), 32); // 16 SGPRs + 16 VGPRs
    }

    #[test]
    fn test_x86_reg_class_display() {
        assert_eq!(X86RegClass::GPR.to_string(), "GPR");
        assert_eq!(X86RegClass::XMM.to_string(), "XMM");
        assert_eq!(X86RegClass::YMM.to_string(), "YMM");
        assert_eq!(X86RegClass::ZMM.to_string(), "ZMM");
        assert_eq!(X86RegClass::Mask.to_string(), "MASK");
    }

    #[test]
    fn test_addr_space_display() {
        assert_eq!(AmdgpuAddrSpace::Global.to_string(), "global");
        assert_eq!(AmdgpuAddrSpace::Local.to_string(), "local");
    }

    #[test]
    fn test_cross_target_allocated_count() {
        let mut ct = CrossTargetAmdgpu::new(AmdgpuIsaVersion::GFX900);
        assert_eq!(ct.allocated_count(), 0);
        ct.allocate_vgpr();
        ct.allocate_vgpr();
        ct.allocate_sgpr();
        assert_eq!(ct.allocated_count(), 3);
    }
}