use std::collections::{BTreeMap, HashMap, HashSet};
use std::fmt;
use super::nvptx_instr_info::{NvptxInstrDesc, NvptxInstrInfo, NvptxOpcode, NvptxOperandType};
use super::nvptx_register_info::{NvptxRegClass, NvptxRegisterInfo};
use super::nvptx_target_machine::NvptxTargetMachine;
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, PartialOrd, Ord)]
pub enum NvptxX86Reg {
RAX,
RBX,
RCX,
RDX,
RSI,
RDI,
RBP,
RSP,
R8,
R9,
R10,
R11,
R12,
R13,
R14,
R15,
XMM0,
XMM1,
XMM2,
XMM3,
XMM4,
XMM5,
XMM6,
XMM7,
XMM8,
XMM9,
XMM10,
XMM11,
XMM12,
XMM13,
XMM14,
XMM15,
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,
}
impl NvptxX86Reg {
pub fn name(&self) -> &'static str {
use NvptxX86Reg::*;
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",
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",
}
}
pub fn reg_class(&self) -> NvptxX86RegClass {
use NvptxX86Reg::*;
match self {
RAX | RBX | RCX | RDX | RSI | RDI | RBP | RSP | R8 | R9 | R10 | R11 | R12 | R13
| R14 | R15 => NvptxX86RegClass::GPR,
XMM0 | XMM1 | XMM2 | XMM3 | XMM4 | XMM5 | XMM6 | XMM7 | XMM8 | XMM9 | XMM10 | XMM11
| XMM12 | XMM13 | XMM14 | XMM15 => NvptxX86RegClass::XMM,
YMM0 | YMM1 | YMM2 | YMM3 | YMM4 | YMM5 | YMM6 | YMM7 | YMM8 | YMM9 | YMM10 | YMM11
| YMM12 | YMM13 | YMM14 | YMM15 => NvptxX86RegClass::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 => {
NvptxX86RegClass::ZMM
}
}
}
pub fn size_bits(&self) -> u32 {
use NvptxX86Reg::*;
match self {
RAX | RBX | RCX | RDX | RSI | RDI | RBP | RSP | R8 | R9 | R10 | R11 | R12 | R13
| R14 | R15 => 64,
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,
}
}
pub fn encoding(&self) -> u32 {
use NvptxX86Reg::*;
match self {
RAX => 0,
RCX => 1,
RDX => 2,
RBX => 3,
RSP => 4,
RBP => 5,
RSI => 6,
RDI => 7,
R8 => 8,
R9 => 9,
R10 => 10,
R11 => 11,
R12 => 12,
R13 => 13,
R14 => 14,
R15 => 15,
_ => 0,
}
}
}
impl fmt::Display for NvptxX86Reg {
fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
write!(f, "{}", self.name())
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum NvptxX86RegClass {
GPR,
XMM,
YMM,
ZMM,
}
impl fmt::Display for NvptxX86RegClass {
fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
match self {
NvptxX86RegClass::GPR => write!(f, "GPR"),
NvptxX86RegClass::XMM => write!(f, "XMM"),
NvptxX86RegClass::YMM => write!(f, "YMM"),
NvptxX86RegClass::ZMM => write!(f, "ZMM"),
}
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum NvptxX86Opcode {
MovRR,
MovRI,
MovMR,
MovRM,
MovZX,
MovSX,
MovSXD,
Push,
Pop,
Xchg,
Lea,
CmovE,
CmovNE,
CmovG,
CmovGE,
CmovL,
CmovLE,
CmovA,
CmovAE,
CmovB,
CmovBE,
Add,
Sub,
Mul,
IMul,
Div,
IDiv,
Neg,
Inc,
Dec,
Adc,
Sbb,
And,
Or,
Xor,
Not,
Shl,
Shr,
Sar,
Rol,
Ror,
Bsf,
Bsr,
Bswap,
Cmp,
Test,
Jmp,
JmpE,
JmpNE,
JmpG,
JmpGE,
JmpL,
JmpLE,
JmpA,
JmpAE,
JmpB,
JmpBE,
Call,
Ret,
Addps,
Subps,
Mulps,
Divps,
Sqrtps,
Rcpps,
Rsqrts,
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,
Minpd,
Maxpd,
Minsd,
Maxsd,
Andps,
Andnps,
Orps,
Xorps,
Cmpps,
Cmppd,
Cmpss,
Cmpsd,
VAddps,
VSubps,
VMulps,
VDivps,
VSqrtps,
VAddpd,
VSubpd,
VMulpd,
VDivpd,
VSqrtpd,
VBroadcastss,
VBroadcastsd,
VPermilps,
VPerm2f128,
VPermq,
VGatherdps,
VGatherdpd,
VGatherqps,
VGatherqpd,
VAddpsZ,
VSubpsZ,
VMulpsZ,
VDivpsZ,
VAddpdZ,
VSubpdZ,
VMulpdZ,
VDivpdZ,
VFmadd132ps,
VFmadd213ps,
VFmadd231ps,
VFmadd132pd,
VFmadd213pd,
VFmadd231pd,
LockAdd,
LockSub,
LockAnd,
LockOr,
LockXor,
LockCmpxchg,
Nop,
Mfence,
Lfence,
Sfence,
CpuId,
Rdtsc,
}
impl NvptxX86Opcode {
pub fn mnemonic(&self) -> &'static str {
use NvptxX86Opcode::*;
match self {
MovRR | MovRI | MovMR | MovRM => "mov",
MovZX => "movzx",
MovSX => "movsx",
MovSXD => "movsxd",
Push => "push",
Pop => "pop",
Xchg => "xchg",
Lea => "lea",
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",
Adc => "adc",
Sbb => "sbb",
And => "and",
Or => "or",
Xor => "xor",
Not => "not",
Shl => "shl",
Shr => "shr",
Sar => "sar",
Rol => "rol",
Ror => "ror",
Bsf => "bsf",
Bsr => "bsr",
Bswap => "bswap",
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",
Rcpps => "rcpps",
Rsqrts => "rsqrtps",
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",
Minpd => "minpd",
Maxpd => "maxpd",
Minsd => "minsd",
Maxsd => "maxsd",
Andps => "andps",
Andnps => "andnps",
Orps => "orps",
Xorps => "xorps",
Cmpps => "cmpps",
Cmppd => "cmppd",
Cmpss => "cmpss",
Cmpsd => "cmpsd",
VAddps => "vaddps",
VSubps => "vsubps",
VMulps => "vmulps",
VDivps => "vdivps",
VSqrtps => "vsqrtps",
VAddpd => "vaddpd",
VSubpd => "vsubpd",
VMulpd => "vmulpd",
VDivpd => "vdivpd",
VSqrtpd => "vsqrtpd",
VBroadcastss => "vbroadcastss",
VBroadcastsd => "vbroadcastsd",
VPermilps => "vpermilps",
VPerm2f128 => "vperm2f128",
VPermq => "vpermq",
VGatherdps => "vgatherdps",
VGatherdpd => "vgatherdpd",
VGatherqps => "vgatherqps",
VGatherqpd => "vgatherqpd",
VAddpsZ => "vaddps",
VSubpsZ => "vsubps",
VMulpsZ => "vmulps",
VDivpsZ => "vdivps",
VAddpdZ => "vaddpd",
VSubpdZ => "vsubpd",
VMulpdZ => "vmulpd",
VDivpdZ => "vdivpd",
VFmadd132ps => "vfmadd132ps",
VFmadd213ps => "vfmadd213ps",
VFmadd231ps => "vfmadd231ps",
VFmadd132pd => "vfmadd132pd",
VFmadd213pd => "vfmadd213pd",
VFmadd231pd => "vfmadd231pd",
LockAdd => "lock add",
LockSub => "lock sub",
LockAnd => "lock and",
LockOr => "lock or",
LockXor => "lock xor",
LockCmpxchg => "lock cmpxchg",
Nop => "nop",
Mfence => "mfence",
Lfence => "lfence",
Sfence => "sfence",
CpuId => "cpuid",
Rdtsc => "rdtsc",
}
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum NvptxSpecialReg {
TidX,
TidY,
TidZ,
NtidX,
NtidY,
NtidZ,
CtaidX,
CtaidY,
CtaidZ,
NctaidX,
NctaidY,
NctaidZ,
LaneId,
WarpId,
SmId,
GridId,
Clock,
Clock64,
Pm0,
Pm1,
Pm2,
Pm3,
Pm4,
Pm5,
Pm6,
Pm7,
EnvReg0,
EnvReg1,
EnvReg2,
EnvReg3,
EnvReg4,
EnvReg5,
GlobalTimerLo,
GlobalTimerHi,
}
impl NvptxSpecialReg {
pub fn ptx_name(&self) -> &'static str {
match self {
Self::TidX => "%tid.x",
Self::TidY => "%tid.y",
Self::TidZ => "%tid.z",
Self::NtidX => "%ntid.x",
Self::NtidY => "%ntid.y",
Self::NtidZ => "%ntid.z",
Self::CtaidX => "%ctaid.x",
Self::CtaidY => "%ctaid.y",
Self::CtaidZ => "%ctaid.z",
Self::NctaidX => "%nctaid.x",
Self::NctaidY => "%nctaid.y",
Self::NctaidZ => "%nctaid.z",
Self::LaneId => "%laneid",
Self::WarpId => "%warpid",
Self::SmId => "%smid",
Self::GridId => "%gridid",
Self::Clock => "%clock",
Self::Clock64 => "%clock64",
Self::Pm0 => "%pm0",
Self::Pm1 => "%pm1",
Self::Pm2 => "%pm2",
Self::Pm3 => "%pm3",
Self::Pm4 => "%pm4",
Self::Pm5 => "%pm5",
Self::Pm6 => "%pm6",
Self::Pm7 => "%pm7",
Self::EnvReg0 => "%envreg0",
Self::EnvReg1 => "%envreg1",
Self::EnvReg2 => "%envreg2",
Self::EnvReg3 => "%envreg3",
Self::EnvReg4 => "%envreg4",
Self::EnvReg5 => "%envreg5",
Self::GlobalTimerLo => "%globaltimer_lo",
Self::GlobalTimerHi => "%globaltimer_hi",
}
}
pub fn is_per_thread(&self) -> bool {
matches!(self, Self::TidX | Self::TidY | Self::TidZ | Self::LaneId)
}
pub fn is_per_block(&self) -> bool {
matches!(
self,
Self::NtidX
| Self::NtidY
| Self::NtidZ
| Self::CtaidX
| Self::CtaidY
| Self::CtaidZ
| Self::NctaidX
| Self::NctaidY
| Self::NctaidZ
| Self::WarpId
| Self::SmId
)
}
pub fn is_per_grid(&self) -> bool {
matches!(self, Self::GridId)
}
pub fn map_to_x86_runtime(&self) -> Option<NvptxX86Reg> {
match self {
Self::TidX => Some(NvptxX86Reg::RAX),
Self::TidY => Some(NvptxX86Reg::RBX),
Self::TidZ => Some(NvptxX86Reg::RCX),
Self::CtaidX => Some(NvptxX86Reg::R8),
Self::CtaidY => Some(NvptxX86Reg::R9),
Self::CtaidZ => Some(NvptxX86Reg::R10),
Self::LaneId => Some(NvptxX86Reg::R11),
Self::WarpId => Some(NvptxX86Reg::R12),
Self::NtidX => Some(NvptxX86Reg::R13),
_ => None,
}
}
}
impl fmt::Display for NvptxSpecialReg {
fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
write!(f, "{}", self.ptx_name())
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum NvptxCallingConvention {
Kernel,
DeviceFunction,
ExternDevice,
Indirect,
Fast,
}
impl NvptxCallingConvention {
pub fn ptx_directive(&self) -> &'static str {
match self {
Self::Kernel => ".entry",
Self::DeviceFunction => ".func",
Self::ExternDevice => ".extern .func",
Self::Indirect => ".func",
Self::Fast => ".func",
}
}
pub fn name(&self) -> &'static str {
match self {
Self::Kernel => "nvptx_kernel",
Self::DeviceFunction => "nvptx_device",
Self::ExternDevice => "nvptx_extern_device",
Self::Indirect => "nvptx_indirect",
Self::Fast => "nvptx_fast",
}
}
pub fn param_space(&self) -> NvptxParamSpace {
match self {
Self::Kernel => NvptxParamSpace::KernelParams,
Self::DeviceFunction | Self::ExternDevice | Self::Fast => NvptxParamSpace::Registers,
Self::Indirect => NvptxParamSpace::Stack,
}
}
pub fn uses_param_space(&self) -> bool {
matches!(self, Self::Kernel)
}
pub fn uses_stack_args(&self) -> bool {
matches!(self, Self::Indirect)
}
pub fn supports_return(&self) -> bool {
!matches!(self, Self::Kernel)
}
pub fn implicit_param_count(&self) -> u32 {
match self {
Self::Kernel => 0,
_ => 0,
}
}
}
impl fmt::Display for NvptxCallingConvention {
fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
write!(f, "{}", self.name())
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum NvptxParamSpace {
KernelParams,
Registers,
Stack,
}
impl fmt::Display for NvptxParamSpace {
fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
match self {
NvptxParamSpace::KernelParams => write!(f, "kernel_params"),
NvptxParamSpace::Registers => write!(f, "registers"),
NvptxParamSpace::Stack => write!(f, "stack"),
}
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub enum NvptxAddrSpace {
Reg,
SpecialReg,
Const,
Global,
Local,
Param,
Shared,
Texture,
Surface,
Generic,
}
impl NvptxAddrSpace {
pub fn ptx_modifier(&self) -> &'static str {
match self {
Self::Reg => ".reg",
Self::SpecialReg => ".sreg",
Self::Const => ".const",
Self::Global => ".global",
Self::Local => ".local",
Self::Param => ".param",
Self::Shared => ".shared",
Self::Texture => ".tex",
Self::Surface => ".surf",
Self::Generic => "",
}
}
pub fn is_writable(&self) -> bool {
match self {
Self::Reg | Self::Global | Self::Local | Self::Shared | Self::Generic => true,
Self::Const | Self::Texture | Self::Param | Self::SpecialReg | Self::Surface => false,
}
}
pub fn is_cached(&self) -> bool {
match self {
Self::Const | Self::Texture | Self::Shared => true,
Self::Global => true,
_ => false,
}
}
pub fn to_llvm_addrspace(&self) -> u32 {
match self {
Self::Reg | Self::SpecialReg => 0,
Self::Global => 1,
Self::Const => 4,
Self::Local => 5,
Self::Param => 101,
Self::Shared => 3,
Self::Texture => 2,
Self::Surface => 2,
Self::Generic => 0,
}
}
pub fn from_llvm_addrspace(n: u32) -> Option<Self> {
match n {
0 => Some(Self::Generic),
1 => Some(Self::Global),
2 => Some(Self::Texture),
3 => Some(Self::Shared),
4 => Some(Self::Const),
5 => Some(Self::Local),
101 => Some(Self::Param),
_ => None,
}
}
}
impl fmt::Display for NvptxAddrSpace {
fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
write!(f, "{}", self.ptx_modifier())
}
}
#[derive(Debug, Clone)]
pub struct NvptxInstructionMapping {
pub ptx_opcode: NvptxOpcode,
pub x86_opcodes: Vec<NvptxX86Opcode>,
pub expansion_factor: u32,
pub simt_aware: bool,
pub is_direct: bool,
pub preserves_semantics: bool,
pub notes: &'static str,
}
impl NvptxInstructionMapping {
pub fn new(
ptx_opcode: NvptxOpcode,
x86_opcodes: Vec<NvptxX86Opcode>,
expansion_factor: u32,
) -> Self {
Self {
ptx_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
}
}
#[derive(Debug, Clone)]
pub struct NvptxRegisterInfoX86 {
pub base: NvptxRegisterInfo,
pub virtual_to_x86: HashMap<u32, NvptxX86Reg>,
pub special_to_x86: HashMap<NvptxSpecialReg, NvptxX86Reg>,
pub max_virtual_reg: u32,
pub use_avx: bool,
}
impl NvptxRegisterInfoX86 {
pub fn new() -> Self {
let base = NvptxRegisterInfo;
let mut special_to_x86 = HashMap::new();
special_to_x86.insert(NvptxSpecialReg::TidX, NvptxX86Reg::RAX);
special_to_x86.insert(NvptxSpecialReg::TidY, NvptxX86Reg::RBX);
special_to_x86.insert(NvptxSpecialReg::TidZ, NvptxX86Reg::RCX);
special_to_x86.insert(NvptxSpecialReg::NtidX, NvptxX86Reg::RDX);
special_to_x86.insert(NvptxSpecialReg::CtaidX, NvptxX86Reg::R8);
special_to_x86.insert(NvptxSpecialReg::CtaidY, NvptxX86Reg::R9);
special_to_x86.insert(NvptxSpecialReg::CtaidZ, NvptxX86Reg::R10);
special_to_x86.insert(NvptxSpecialReg::LaneId, NvptxX86Reg::R11);
special_to_x86.insert(NvptxSpecialReg::WarpId, NvptxX86Reg::R12);
special_to_x86.insert(NvptxSpecialReg::SmId, NvptxX86Reg::R13);
Self {
base,
virtual_to_x86: HashMap::new(),
special_to_x86,
max_virtual_reg: 0,
use_avx: true,
}
}
pub fn allocate_virtual_reg(&mut self, width: u32) -> (u32, NvptxX86Reg) {
let reg_id = self.max_virtual_reg;
self.max_virtual_reg += 1;
let x86_reg = if width <= 64 {
let gpr_id = reg_id % 12;
match gpr_id {
0 => NvptxX86Reg::RAX,
1 => NvptxX86Reg::RBX,
2 => NvptxX86Reg::RCX,
3 => NvptxX86Reg::RDX,
4 => NvptxX86Reg::RSI,
5 => NvptxX86Reg::RDI,
6 => NvptxX86Reg::R8,
7 => NvptxX86Reg::R9,
8 => NvptxX86Reg::R10,
9 => NvptxX86Reg::R11,
10 => NvptxX86Reg::R12,
11 => NvptxX86Reg::R13,
_ => NvptxX86Reg::R14,
}
} else if width <= 128 || !self.use_avx {
let xmm_id = reg_id % 16;
match xmm_id {
0 => NvptxX86Reg::XMM0,
1 => NvptxX86Reg::XMM1,
2 => NvptxX86Reg::XMM2,
3 => NvptxX86Reg::XMM3,
4 => NvptxX86Reg::XMM4,
5 => NvptxX86Reg::XMM5,
6 => NvptxX86Reg::XMM6,
7 => NvptxX86Reg::XMM7,
8 => NvptxX86Reg::XMM8,
9 => NvptxX86Reg::XMM9,
10 => NvptxX86Reg::XMM10,
11 => NvptxX86Reg::XMM11,
12 => NvptxX86Reg::XMM12,
13 => NvptxX86Reg::XMM13,
14 => NvptxX86Reg::XMM14,
15 => NvptxX86Reg::XMM15,
_ => unreachable!(),
}
} else if width <= 256 {
let ymm_id = reg_id % 16;
match ymm_id {
0 => NvptxX86Reg::YMM0,
1 => NvptxX86Reg::YMM1,
2 => NvptxX86Reg::YMM2,
3 => NvptxX86Reg::YMM3,
4 => NvptxX86Reg::YMM4,
5 => NvptxX86Reg::YMM5,
6 => NvptxX86Reg::YMM6,
7 => NvptxX86Reg::YMM7,
8 => NvptxX86Reg::YMM8,
9 => NvptxX86Reg::YMM9,
10 => NvptxX86Reg::YMM10,
11 => NvptxX86Reg::YMM11,
12 => NvptxX86Reg::YMM12,
13 => NvptxX86Reg::YMM13,
14 => NvptxX86Reg::YMM14,
15 => NvptxX86Reg::YMM15,
_ => unreachable!(),
}
} else {
let zmm_id = reg_id % 16;
match zmm_id {
0 => NvptxX86Reg::ZMM0,
1 => NvptxX86Reg::ZMM1,
2 => NvptxX86Reg::ZMM2,
3 => NvptxX86Reg::ZMM3,
4 => NvptxX86Reg::ZMM4,
5 => NvptxX86Reg::ZMM5,
6 => NvptxX86Reg::ZMM6,
7 => NvptxX86Reg::ZMM7,
8 => NvptxX86Reg::ZMM8,
9 => NvptxX86Reg::ZMM9,
10 => NvptxX86Reg::ZMM10,
11 => NvptxX86Reg::ZMM11,
12 => NvptxX86Reg::ZMM12,
13 => NvptxX86Reg::ZMM13,
14 => NvptxX86Reg::ZMM14,
15 => NvptxX86Reg::ZMM15,
_ => unreachable!(),
}
};
self.virtual_to_x86.insert(reg_id, x86_reg);
(reg_id, x86_reg)
}
pub fn free_virtual_reg(&mut self, reg_id: u32) {
self.virtual_to_x86.remove(®_id);
}
pub fn lookup_x86(&self, virtual_id: u32) -> Option<NvptxX86Reg> {
self.virtual_to_x86.get(&virtual_id).copied()
}
pub fn lookup_special_x86(&self, special: NvptxSpecialReg) -> Option<NvptxX86Reg> {
self.special_to_x86.get(&special).copied()
}
pub fn allocated_virtual_count(&self) -> usize {
self.virtual_to_x86.len()
}
pub fn mapped_special_count(&self) -> usize {
self.special_to_x86.len()
}
pub fn reset(&mut self) {
self.virtual_to_x86.clear();
self.max_virtual_reg = 0;
}
}
#[derive(Debug, Clone)]
pub struct NvptxX86TargetMachine {
pub base: NvptxTargetMachine,
pub x86_register_info: NvptxRegisterInfoX86,
pub instruction_mappings: Vec<NvptxInstructionMapping>,
pub x86_triple: String,
pub x86_data_layout: String,
pub opt_level: u32,
pub use_avx: bool,
pub use_avx2: bool,
pub use_avx512: bool,
pub emit_simt_masks: bool,
pub sm_version: u32,
}
impl NvptxX86TargetMachine {
pub fn new(sm_version: u32) -> Self {
let base = NvptxTargetMachine::new(sm_version, true);
let x86_register_info = NvptxRegisterInfoX86::new();
let instruction_mappings = build_ptx_x86_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-".to_string()
+ "f80:128-n8:16:32:64-S128",
opt_level: 2,
use_avx: true,
use_avx2: true,
use_avx512: false,
emit_simt_masks: true,
sm_version,
}
}
pub fn set_opt_level(&mut self, level: u32) {
self.opt_level = level.min(3);
}
pub fn enable_avx512(&mut self) {
self.use_avx512 = true;
}
pub fn x86_feature_string(&self) -> String {
let mut features = vec!["+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.join(",")
}
pub fn lookup_mapping(&self, opcode: NvptxOpcode) -> Option<&NvptxInstructionMapping> {
self.instruction_mappings
.iter()
.find(|m| m.ptx_opcode == opcode)
}
pub fn all_mappings(&self) -> &[NvptxInstructionMapping] {
&self.instruction_mappings
}
pub fn mapping_count(&self) -> usize {
self.instruction_mappings.len()
}
pub fn get_x86_data_layout(&self) -> &str {
&self.x86_data_layout
}
pub fn compute_capability(&self) -> f32 {
(self.sm_version as f32) / 10.0
}
}
fn build_ptx_x86_mappings() -> Vec<NvptxInstructionMapping> {
let mut m = Vec::new();
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Add,
vec![NvptxX86Opcode::Add],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Sub,
vec![NvptxX86Opcode::Sub],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Mul,
vec![NvptxX86Opcode::IMul],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Mad,
vec![NvptxX86Opcode::IMul, NvptxX86Opcode::Add],
2,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Div,
vec![NvptxX86Opcode::IDiv],
1,
));
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::Rem,
vec![NvptxX86Opcode::IDiv, NvptxX86Opcode::MovRR],
2,
)
.with_notes("rem via div + mov remainder"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Abs,
vec![
NvptxX86Opcode::MovRR,
NvptxX86Opcode::Neg,
NvptxX86Opcode::CmovL,
],
3,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Neg,
vec![NvptxX86Opcode::Neg],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Min,
vec![NvptxX86Opcode::Cmp, NvptxX86Opcode::CmovL],
2,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Max,
vec![NvptxX86Opcode::Cmp, NvptxX86Opcode::CmovG],
2,
));
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::Popc,
vec![
NvptxX86Opcode::MovRR,
NvptxX86Opcode::Shr,
NvptxX86Opcode::And,
NvptxX86Opcode::Add,
],
4,
)
.with_notes("popcount via software sequence"),
);
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::Clz,
vec![NvptxX86Opcode::Bsr, NvptxX86Opcode::Xor],
2,
)
.with_notes("clz via bsr + xor if using LZCNT"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Bfind,
vec![NvptxX86Opcode::Bsf],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Brev,
vec![NvptxX86Opcode::Bswap],
1,
));
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::Sad,
vec![
NvptxX86Opcode::Sub,
NvptxX86Opcode::Abs,
NvptxX86Opcode::Add,
],
3,
)
.with_notes("SAD via absdiff + accumulate"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Bfe,
vec![NvptxX86Opcode::Shr, NvptxX86Opcode::And],
2,
));
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::Bfi,
vec![
NvptxX86Opcode::And,
NvptxX86Opcode::Shl,
NvptxX86Opcode::And,
NvptxX86Opcode::Or,
],
4,
)
.with_notes("bitfield insert"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::And,
vec![NvptxX86Opcode::And],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Or,
vec![NvptxX86Opcode::Or],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Xor,
vec![NvptxX86Opcode::Xor],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Not,
vec![NvptxX86Opcode::Not],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Shl,
vec![NvptxX86Opcode::Shl],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Shr,
vec![NvptxX86Opcode::Shr],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Cnot,
vec![NvptxX86Opcode::Xor, NvptxX86Opcode::Not],
2,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Mov,
vec![NvptxX86Opcode::MovRR],
1,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Cvt, vec![NvptxX86Opcode::MovRR], 1)
.with_notes("cvt is type-polymorphic; mapping depends on src/dst types"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Cvta, vec![NvptxX86Opcode::Lea], 1)
.with_notes("cvta.to.global maps to LEA on X86"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Ld,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::St,
vec![NvptxX86Opcode::MovMR],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::LdGlobal,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::StGlobal,
vec![NvptxX86Opcode::MovMR],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::LdShared,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::StShared,
vec![NvptxX86Opcode::MovMR],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::LdLocal,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::StLocal,
vec![NvptxX86Opcode::MovMR],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::LdParam,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::StParam,
vec![NvptxX86Opcode::MovMR],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::LdConst,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::LdGlobalNc, vec![NvptxX86Opcode::MovRM], 1)
.with_notes("non-coherent load = plain MOV"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::LdGlobalCg,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::LdGlobalCa,
vec![NvptxX86Opcode::MovRM],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpEq,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpNe,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpLt,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpLe,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpGt,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpGe,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpLo,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpLs,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpHi,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpHs,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpEqu,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpNeu,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpLtu,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpLeu,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpGtu,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpGeu,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpNum,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::SetpNan,
vec![NvptxX86Opcode::Cmp],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Selp,
vec![NvptxX86Opcode::CmovNE],
1,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Bra, vec![NvptxX86Opcode::Jmp], 1)
.with_simt(true),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Call,
vec![NvptxX86Opcode::Call],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Ret,
vec![NvptxX86Opcode::Ret],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Exit,
vec![NvptxX86Opcode::Ret],
1,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Brkpt, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::FAdd,
vec![NvptxX86Opcode::Addss],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::FSub,
vec![NvptxX86Opcode::Subss],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::FMul,
vec![NvptxX86Opcode::Mulss],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::FDiv,
vec![NvptxX86Opcode::Divss],
1,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Fma, vec![NvptxX86Opcode::VFmadd132ps], 1)
.with_notes("FMA → vfmadd on AVX, expands to mul+add otherwise"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::FMin,
vec![NvptxX86Opcode::Minss],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::FMax,
vec![NvptxX86Opcode::Maxss],
1,
));
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::Rcp,
vec![NvptxX86Opcode::MovRI, NvptxX86Opcode::Divss],
2,
)
.with_notes("rcp = 1.0/x, approximate on GPU"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Sqrt,
vec![NvptxX86Opcode::Sqrtss],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Rsqrt,
vec![
NvptxX86Opcode::Sqrtss,
NvptxX86Opcode::MovRI,
NvptxX86Opcode::Divss,
],
3,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Sin, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Cos, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Lg2, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Ex2, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::FAbs, vec![NvptxX86Opcode::Andps], 1)
.with_notes("FAbs via ANDPS with sign-mask clear"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::FNeg, vec![NvptxX86Opcode::Xorps], 1)
.with_notes("FNeg via XORPS with sign-bit toggle"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::FMad,
vec![NvptxX86Opcode::Mulss, NvptxX86Opcode::Addss],
2,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Tanh, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Cvt, vec![NvptxX86Opcode::MovRR], 1)
.with_notes("cvt is type-polymorphic; mapping depends on src/dst types"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::CvtRp, vec![NvptxX86Opcode::Cvttss2si], 1)
.with_notes("cvt.rp (round toward +inf)"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::CvtRn, vec![NvptxX86Opcode::Cvttss2si], 1)
.with_notes("cvt.rn (round to nearest even)"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::CvtRz, vec![NvptxX86Opcode::Cvttss2si], 1)
.with_notes("cvt.rz (round toward zero)"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::CvtRm, vec![NvptxX86Opcode::Cvtss2si], 1)
.with_notes("cvt.rm (round toward -inf)"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomAdd,
vec![NvptxX86Opcode::LockAdd],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomSub,
vec![NvptxX86Opcode::LockSub],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomExch,
vec![NvptxX86Opcode::Xchg],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomAnd,
vec![NvptxX86Opcode::LockAnd],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomOr,
vec![NvptxX86Opcode::LockOr],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomXor,
vec![NvptxX86Opcode::LockXor],
1,
));
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomCas,
vec![NvptxX86Opcode::LockCmpxchg],
1,
));
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::AtomMin,
vec![
NvptxX86Opcode::MovRM,
NvptxX86Opcode::Cmp,
NvptxX86Opcode::CmovL,
NvptxX86Opcode::MovMR,
],
4,
)
.with_notes("min via cmpxchg loop"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::AtomMax,
vec![
NvptxX86Opcode::MovRM,
NvptxX86Opcode::Cmp,
NvptxX86Opcode::CmovG,
NvptxX86Opcode::MovMR,
],
4,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::RedAdd, vec![NvptxX86Opcode::LockAdd], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::RedMin,
vec![
NvptxX86Opcode::MovRM,
NvptxX86Opcode::Cmp,
NvptxX86Opcode::CmovL,
NvptxX86Opcode::MovMR,
],
4,
)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::RedMax,
vec![
NvptxX86Opcode::MovRM,
NvptxX86Opcode::Cmp,
NvptxX86Opcode::CmovG,
NvptxX86Opcode::MovMR,
],
4,
)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::RedAnd, vec![NvptxX86Opcode::LockAnd], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::RedOr, vec![NvptxX86Opcode::LockOr], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::RedSub, vec![NvptxX86Opcode::LockSub], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::RedXor, vec![NvptxX86Opcode::LockXor], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Bar, vec![NvptxX86Opcode::Mfence], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::BarSync, vec![NvptxX86Opcode::Mfence], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::BarArrive, vec![NvptxX86Opcode::Sfence], 1)
.with_simt(true)
.with_notes("bar.arrive → sfence (release semantics)"),
);
m.push(NvptxInstructionMapping::new(
NvptxOpcode::Membar,
vec![NvptxX86Opcode::Mfence],
1,
));
m.push(
NvptxInstructionMapping::new(NvptxOpcode::VoteAll, vec![NvptxX86Opcode::And], 1)
.with_simt(true)
.with_notes("vote.all → reduce across lanes via AND"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::VoteAny, vec![NvptxX86Opcode::Or], 1)
.with_simt(true)
.with_notes("vote.any → reduce across lanes via OR"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::VoteUni, vec![NvptxX86Opcode::Cmp], 1)
.with_simt(true),
);
m.push(
NvptxInstructionMapping::new(
NvptxOpcode::VoteBallot,
vec![NvptxX86Opcode::Shl, NvptxX86Opcode::Or],
2,
)
.with_simt(true)
.with_notes("ballot → build bitmask"),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Tex, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Tld4, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Txq, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Suld, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Sust, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Sured, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Suq, vec![NvptxX86Opcode::Call], 1)
.with_semantics(false),
);
m.push(
NvptxInstructionMapping::new(NvptxOpcode::Isspacep, vec![NvptxX86Opcode::Cmp], 1)
.with_notes("isspacep → check address space via comparison"),
);
m
}
#[derive(Debug, Clone)]
pub struct NVPTXX86Bridge {
pub target_machine: NvptxX86TargetMachine,
pub calling_convention: NvptxCallingConvention,
pub nvptx_instr_info: NvptxInstrInfo,
pub cross_target: CrossTargetNvptx,
pub stats: NvptxBridgeStats,
pub kernel_params: Vec<NvptxKernelParam>,
pub is_64bit: bool,
}
impl NVPTXX86Bridge {
pub fn new(sm_version: u32) -> Self {
let target_machine = NvptxX86TargetMachine::new(sm_version);
let nvptx_instr_info = NvptxInstrInfo::new();
let cross_target = CrossTargetNvptx::new(sm_version);
Self {
target_machine,
calling_convention: NvptxCallingConvention::Kernel,
nvptx_instr_info,
cross_target,
stats: NvptxBridgeStats::default(),
kernel_params: Vec::new(),
is_64bit: true,
}
}
pub fn set_calling_convention(&mut self, cc: NvptxCallingConvention) {
self.calling_convention = cc;
}
pub fn set_64bit(&mut self, is_64bit: bool) {
self.is_64bit = is_64bit;
}
pub fn add_kernel_param(&mut self, param: NvptxKernelParam) {
self.kernel_params.push(param);
}
pub fn clear_kernel_params(&mut self) {
self.kernel_params.clear();
}
pub fn translate_instruction(
&self,
ptx_opcode: NvptxOpcode,
) -> Option<(Vec<NvptxX86Opcode>, u32)> {
self.target_machine
.lookup_mapping(ptx_opcode)
.map(|m| (m.x86_opcodes.clone(), m.expansion_factor))
}
pub fn has_direct_mapping(&self, opcode: NvptxOpcode) -> bool {
self.target_machine
.lookup_mapping(opcode)
.map(|m| m.is_direct)
.unwrap_or(false)
}
pub fn compute_expansion(&self, opcodes: &[NvptxOpcode]) -> u32 {
opcodes
.iter()
.filter_map(|op| self.target_machine.lookup_mapping(*op))
.map(|m| m.expansion_factor)
.sum()
}
pub fn sm_version(&self) -> u32 {
self.target_machine.sm_version
}
pub fn compute_capability(&self) -> f32 {
self.target_machine.compute_capability()
}
pub fn enable_avx512(&mut self) {
self.target_machine.enable_avx512();
self.stats.avx512_enabled = true;
}
pub fn set_opt_level(&mut self, level: u32) {
self.target_machine.set_opt_level(level);
}
pub fn x86_feature_string(&self) -> String {
self.target_machine.x86_feature_string()
}
pub fn get_x86_data_layout(&self) -> &str {
self.target_machine.get_x86_data_layout()
}
pub fn get_nvptx_triple(&self) -> String {
if self.is_64bit {
"nvptx64-nvidia-cuda".to_string()
} else {
"nvptx-nvidia-cuda".to_string()
}
}
pub fn emit_x86_prologue(&self) -> String {
let mut buf = String::new();
buf.push_str("; -- NVPTX Kernel Prologue (X86 Bridge) --\n");
buf.push_str(&format!(
"; PTX Target: nvptx{}-nvidia-cuda\n",
if self.is_64bit { "64" } else { "" }
));
buf.push_str(&format!("; SM Version: sm_{}\n", self.sm_version()));
buf.push_str(&format!(
"; Compute Capability: {:.1}\n",
self.compute_capability()
));
buf.push_str(&format!(
"; Calling Convention: {}\n",
self.calling_convention
));
if self.calling_convention == NvptxCallingConvention::Kernel {
buf.push_str("; Kernel Parameters:\n");
buf.push_str(&format!("; Total params: {}\n", self.kernel_params.len()));
for param in &self.kernel_params {
buf.push_str(&format!(
"; param {} (size {}, align {}, space: {})\n",
param.name, param.size_bytes, param.alignment, param.addr_space
));
}
buf.push_str("; Special Registers: %tid, %ntid, %ctaid, %nctaid, %laneid\n");
}
buf
}
pub fn emit_x86_epilogue(&self) -> String {
let mut buf = String::new();
buf.push_str("; -- NVPTX Kernel Epilogue (X86 Bridge) --\n");
buf.push_str("; ret\n");
buf
}
pub fn allocate_virtual_reg(&mut self, width: u32) -> (u32, NvptxX86Reg) {
self.target_machine
.x86_register_info
.allocate_virtual_reg(width)
}
pub fn get_stats(&self) -> &NvptxBridgeStats {
&self.stats
}
pub fn reset_stats(&mut self) {
self.stats = NvptxBridgeStats::default();
}
pub fn is_ready(&self) -> bool {
self.target_machine.mapping_count() > 0
}
}
#[derive(Debug, Clone, Default)]
pub struct NvptxBridgeStats {
pub instructions_translated: u64,
pub kernels_processed: u64,
pub total_expansion: u64,
pub simt_translations: u64,
pub avx512_enabled: bool,
pub virtual_regs_allocated: u64,
}
#[derive(Debug, Clone)]
pub struct NvptxKernelParam {
pub name: String,
pub size_bytes: u32,
pub alignment: u32,
pub offset: u32,
pub is_pointer: bool,
pub is_byval: bool,
pub addr_space: NvptxAddrSpace,
}
impl NvptxKernelParam {
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: NvptxAddrSpace::Param,
}
}
pub fn with_pointer(mut self, is_ptr: bool) -> Self {
self.is_pointer = is_ptr;
self
}
pub fn with_addr_space(mut self, space: NvptxAddrSpace) -> Self {
self.addr_space = space;
self
}
}
#[derive(Debug, Clone)]
pub struct CrossTargetNvptx {
pub sm_version: u32,
pub warp_size: u32,
pub max_threads_per_block: u32,
pub scalarize_uniform: bool,
pub emit_lane_masks: bool,
pub exec_mask_reg: Option<NvptxX86Reg>,
pub lane_id_reg: Option<NvptxX86Reg>,
pub active_warps: u32,
}
impl CrossTargetNvptx {
pub fn new(sm_version: u32) -> Self {
Self {
sm_version,
warp_size: 32,
max_threads_per_block: 1024,
scalarize_uniform: true,
emit_lane_masks: true,
exec_mask_reg: None,
lane_id_reg: None,
active_warps: 0,
}
}
pub fn set_exec_mask_reg(&mut self, reg: NvptxX86Reg) {
self.exec_mask_reg = Some(reg);
}
pub fn set_lane_id_reg(&mut self, reg: NvptxX86Reg) {
self.lane_id_reg = Some(reg);
}
pub fn emit_lane_init(&self) -> String {
let mut buf = String::new();
buf.push_str("; -- PTX SIMT Lane Initialization (X86 Bridge) --\n");
buf.push_str(&format!("; Warp size: {}\n", self.warp_size));
buf.push_str(&format!(
"; Max threads/block: {}\n",
self.max_threads_per_block
));
buf.push_str("; On X86, each thread is emulated by a slice of a vector register\n");
buf.push_str("; lane_id = thread_idx & 0x1F (lower 5 bits)\n");
buf.push_str("; warp_id = thread_idx >> 5\n");
buf
}
pub fn emit_exec_predicate(&self, inverted: bool) -> String {
if !self.emit_lane_masks {
return String::new();
}
if inverted {
"; bt [exec_mask], eax ; jnc .L_skip\n".to_string()
} else {
"; bt [exec_mask], eax ; jc .L_exec\n".to_string()
}
}
pub fn full_warp_mask() -> u32 {
0xFFFF_FFFF
}
pub fn warps_per_block(&self, block_size: u32) -> u32 {
(block_size + self.warp_size - 1) / self.warp_size
}
pub fn threads_from_warps(&self, num_warps: u32) -> u32 {
num_warps * self.warp_size
}
pub fn emit_warp_loop_header(&self) -> String {
let mut buf = String::new();
buf.push_str("; -- Warp Loop Header (X86) --\n");
buf.push_str("; for (warp = 0; warp < num_warps; warp++) {\n");
buf.push_str("; active_mask = get_active_lanes(warp);\n");
buf.push_str("; for (lane = 0; lane < 32; lane++) {\n");
buf.push_str("; if (active_mask & (1 << lane)) {\n");
buf.push_str("; ... execute lane body ...\n");
buf.push_str("; }\n");
buf.push_str("; }\n");
buf.push_str("; }\n");
buf
}
pub fn is_uniform_instruction(opcode: NvptxOpcode) -> bool {
matches!(
opcode,
NvptxOpcode::Bar | NvptxOpcode::BarSync | NvptxOpcode::BarArrive | NvptxOpcode::Membar
)
}
pub fn reset(&mut self) {
self.exec_mask_reg = None;
self.lane_id_reg = None;
self.active_warps = 0;
}
}
#[derive(Debug, Clone)]
pub struct NvptxKernelDescriptorX86 {
pub name: String,
pub params: Vec<NvptxKernelParam>,
pub register_count: u32,
pub shared_memory: u32,
pub max_threads_per_block: u32,
pub min_threads_per_block: u32,
pub min_sm_version: u32,
pub uses_dynamic_shared: bool,
}
impl NvptxKernelDescriptorX86 {
pub fn new(name: &str) -> Self {
Self {
name: name.to_string(),
params: Vec::new(),
register_count: 0,
shared_memory: 0,
max_threads_per_block: 1024,
min_threads_per_block: 1,
min_sm_version: 50,
uses_dynamic_shared: false,
}
}
pub fn add_param(&mut self, param: NvptxKernelParam) {
self.params.push(param);
}
pub fn set_register_usage(&mut self, count: u32) {
self.register_count = count;
}
pub fn set_shared_memory(&mut self, bytes: u32, dynamic: bool) {
self.shared_memory = bytes;
self.uses_dynamic_shared = dynamic;
}
pub fn emit_x86_launch_stub(&self) -> String {
let mut buf = String::new();
buf.push_str(&format!(
"// X86 Launch Stub for CUDA Kernel: {}\n",
self.name
));
buf.push_str("// This stub would be called by the CUDA runtime on X86\n");
buf.push_str(&format!(
"// cudaLaunchKernel(\"{}\", grid_dim, block_dim, args, shared_mem={});\n",
self.name, self.shared_memory
));
buf.push_str(&format!(
"// reg_count={}, max_threads_per_block={}\n",
self.register_count, self.max_threads_per_block
));
buf
}
pub fn emit_ptx_entry(&self) -> String {
let mut buf = String::new();
buf.push_str(&format!(".visible .entry {} (\n", self.name));
for (i, param) in self.params.iter().enumerate() {
let comma = if i < self.params.len() - 1 { "," } else { "" };
buf.push_str(&format!(
" .param .u{} param_{}{}\n",
param.size_bytes * 8,
i,
comma
));
}
buf.push_str(")\n");
buf.push_str(&format!(".maxntid {}, 1, 1\n", self.max_threads_per_block));
buf.push_str(".minnctapersm 4\n");
buf.push_str(&format!(".maxnreg {}\n", self.register_count));
buf.push_str("{\n");
buf.push_str(" // Kernel body\n");
buf.push_str(" ret;\n");
buf.push_str("}\n");
buf
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum NvptxMemoryOrder {
Relaxed,
Acquire,
Release,
AcquireRelease,
SequentiallyConsistent,
}
impl NvptxMemoryOrder {
pub fn to_ptx_suffix(&self) -> &'static str {
match self {
Self::Relaxed => ".relaxed",
Self::Acquire => ".acquire",
Self::Release => ".release",
Self::AcquireRelease => ".acq_rel",
Self::SequentiallyConsistent => ".sc",
}
}
pub fn to_x86_barrier(&self) -> Option<&'static str> {
match self {
Self::Relaxed | Self::Acquire | Self::Release => None,
Self::AcquireRelease | Self::SequentiallyConsistent => Some("mfence"),
}
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum NvptxMemoryScope {
Cta,
Gpu,
Sys,
}
impl NvptxMemoryScope {
pub fn to_ptx_suffix(&self) -> &'static str {
match self {
Self::Cta => ".cta",
Self::Gpu => ".gpu",
Self::Sys => ".sys",
}
}
pub fn to_x86_barrier(&self) -> Option<&'static str> {
match self {
Self::Cta => None,
Self::Gpu | Self::Sys => Some("mfence"),
}
}
}
#[derive(Debug, Clone)]
pub struct NvptxWarpEmulator {
pub warp_size: u32,
pub active_mask: u32,
}
impl NvptxWarpEmulator {
pub fn new() -> Self {
Self {
warp_size: 32,
active_mask: 0xFFFF_FFFF,
}
}
pub fn emit_shfl_sync(&self, src_lane: u32, width: u32) -> String {
format!(
"; shfl.sync: broadcast from lane {} within groups of {}\n",
src_lane, width
)
}
pub fn emit_ballot(&self) -> String {
"; ballot: collect predicate masks across warp\n".to_string()
}
pub fn all_active(&self) -> bool {
self.active_mask == 0xFFFF_FFFF
}
pub fn deactivate_lane(&mut self, lane: u32) {
self.active_mask &= !(1 << lane);
}
pub fn activate_lane(&mut self, lane: u32) {
self.active_mask |= 1 << lane;
}
}
#[derive(Debug, Clone)]
pub struct NvptxBridgeBuilder {
sm_version: u32,
is_64bit: bool,
opt_level: u32,
use_avx: bool,
use_avx2: bool,
use_avx512: bool,
calling_convention: NvptxCallingConvention,
emit_simt: bool,
scalarize: bool,
kernel_params: Vec<NvptxKernelParam>,
}
impl NvptxBridgeBuilder {
pub fn new(sm_version: u32) -> Self {
Self {
sm_version,
is_64bit: true,
opt_level: 2,
use_avx: true,
use_avx2: true,
use_avx512: false,
calling_convention: NvptxCallingConvention::Kernel,
emit_simt: true,
scalarize: true,
kernel_params: Vec::new(),
}
}
pub fn bit32(mut self) -> Self {
self.is_64bit = false;
self
}
pub fn bit64(mut self) -> Self {
self.is_64bit = true;
self
}
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: NvptxCallingConvention) -> 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_param(mut self, param: NvptxKernelParam) -> Self {
self.kernel_params.push(param);
self
}
pub fn build(self) -> NVPTXX86Bridge {
let mut bridge = NVPTXX86Bridge::new(self.sm_version);
bridge.set_64bit(self.is_64bit);
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.emit_lane_masks = self.emit_simt;
bridge.cross_target.scalarize_uniform = self.scalarize;
for param in self.kernel_params {
bridge.add_kernel_param(param);
}
bridge
}
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn test_nvptx_x86_reg_names() {
assert_eq!(NvptxX86Reg::RAX.name(), "rax");
assert_eq!(NvptxX86Reg::RBX.name(), "rbx");
assert_eq!(NvptxX86Reg::XMM0.name(), "xmm0");
assert_eq!(NvptxX86Reg::YMM0.name(), "ymm0");
assert_eq!(NvptxX86Reg::ZMM0.name(), "zmm0");
}
#[test]
fn test_nvptx_x86_reg_sizes() {
assert_eq!(NvptxX86Reg::RAX.size_bits(), 64);
assert_eq!(NvptxX86Reg::XMM0.size_bits(), 128);
assert_eq!(NvptxX86Reg::YMM0.size_bits(), 256);
assert_eq!(NvptxX86Reg::ZMM0.size_bits(), 512);
}
#[test]
fn test_nvptx_x86_reg_classes() {
assert_eq!(NvptxX86Reg::RAX.reg_class(), NvptxX86RegClass::GPR);
assert_eq!(NvptxX86Reg::XMM0.reg_class(), NvptxX86RegClass::XMM);
assert_eq!(NvptxX86Reg::YMM0.reg_class(), NvptxX86RegClass::YMM);
assert_eq!(NvptxX86Reg::ZMM0.reg_class(), NvptxX86RegClass::ZMM);
}
#[test]
fn test_nvptx_x86_reg_display() {
assert_eq!(format!("{}", NvptxX86Reg::RAX), "rax");
assert_eq!(format!("{}", NvptxX86Reg::XMM15), "xmm15");
}
#[test]
fn test_nvptx_x86_reg_encoding() {
assert_eq!(NvptxX86Reg::RAX.encoding(), 0);
assert_eq!(NvptxX86Reg::RCX.encoding(), 1);
assert_eq!(NvptxX86Reg::RDX.encoding(), 2);
assert_eq!(NvptxX86Reg::RBX.encoding(), 3);
}
#[test]
fn test_nvptx_x86_opcode_mnemonics() {
assert_eq!(NvptxX86Opcode::Add.mnemonic(), "add");
assert_eq!(NvptxX86Opcode::Sub.mnemonic(), "sub");
assert_eq!(NvptxX86Opcode::IMul.mnemonic(), "imul");
assert_eq!(NvptxX86Opcode::And.mnemonic(), "and");
assert_eq!(NvptxX86Opcode::Or.mnemonic(), "or");
assert_eq!(NvptxX86Opcode::Xor.mnemonic(), "xor");
assert_eq!(NvptxX86Opcode::Nop.mnemonic(), "nop");
assert_eq!(NvptxX86Opcode::Ret.mnemonic(), "ret");
}
#[test]
fn test_nvptx_x86_sse_opcodes() {
assert_eq!(NvptxX86Opcode::Addss.mnemonic(), "addss");
assert_eq!(NvptxX86Opcode::Subsd.mnemonic(), "subsd");
assert_eq!(NvptxX86Opcode::Sqrtps.mnemonic(), "sqrtps");
}
#[test]
fn test_nvptx_x86_avx_opcodes() {
assert_eq!(NvptxX86Opcode::VAddps.mnemonic(), "vaddps");
assert_eq!(NvptxX86Opcode::VBroadcastss.mnemonic(), "vbroadcastss");
}
#[test]
fn test_nvptx_x86_atomic_opcodes() {
assert_eq!(NvptxX86Opcode::LockAdd.mnemonic(), "lock add");
assert_eq!(NvptxX86Opcode::LockCmpxchg.mnemonic(), "lock cmpxchg");
}
#[test]
fn test_special_reg_ptx_names() {
assert_eq!(NvptxSpecialReg::TidX.ptx_name(), "%tid.x");
assert_eq!(NvptxSpecialReg::NtidY.ptx_name(), "%ntid.y");
assert_eq!(NvptxSpecialReg::CtaidZ.ptx_name(), "%ctaid.z");
assert_eq!(NvptxSpecialReg::LaneId.ptx_name(), "%laneid");
assert_eq!(NvptxSpecialReg::WarpId.ptx_name(), "%warpid");
assert_eq!(NvptxSpecialReg::SmId.ptx_name(), "%smid");
}
#[test]
fn test_special_reg_per_thread() {
assert!(NvptxSpecialReg::TidX.is_per_thread());
assert!(NvptxSpecialReg::LaneId.is_per_thread());
assert!(!NvptxSpecialReg::CtaidX.is_per_thread());
}
#[test]
fn test_special_reg_per_block() {
assert!(NvptxSpecialReg::CtaidX.is_per_block());
assert!(NvptxSpecialReg::NtidX.is_per_block());
assert!(NvptxSpecialReg::WarpId.is_per_block());
assert!(!NvptxSpecialReg::LaneId.is_per_block());
}
#[test]
fn test_special_reg_display() {
assert_eq!(NvptxSpecialReg::TidX.to_string(), "%tid.x");
assert_eq!(NvptxSpecialReg::LaneId.to_string(), "%laneid");
}
#[test]
fn test_special_reg_map_to_x86() {
assert_eq!(
NvptxSpecialReg::TidX.map_to_x86_runtime(),
Some(NvptxX86Reg::RAX)
);
assert_eq!(
NvptxSpecialReg::CtaidX.map_to_x86_runtime(),
Some(NvptxX86Reg::R8)
);
assert_eq!(
NvptxSpecialReg::LaneId.map_to_x86_runtime(),
Some(NvptxX86Reg::R11)
);
}
#[test]
fn test_kernel_calling_convention() {
let cc = NvptxCallingConvention::Kernel;
assert_eq!(cc.ptx_directive(), ".entry");
assert!(cc.uses_param_space());
assert!(!cc.uses_stack_args());
assert!(!cc.supports_return());
}
#[test]
fn test_device_calling_convention() {
let cc = NvptxCallingConvention::DeviceFunction;
assert_eq!(cc.ptx_directive(), ".func");
assert!(!cc.uses_param_space());
assert!(!cc.uses_stack_args());
assert!(cc.supports_return());
}
#[test]
fn test_indirect_calling_convention() {
let cc = NvptxCallingConvention::Indirect;
assert!(cc.uses_stack_args());
}
#[test]
fn test_calling_convention_display() {
assert_eq!(NvptxCallingConvention::Kernel.to_string(), "nvptx_kernel");
assert_eq!(
NvptxCallingConvention::DeviceFunction.to_string(),
"nvptx_device"
);
}
#[test]
fn test_param_space_display() {
assert_eq!(NvptxParamSpace::KernelParams.to_string(), "kernel_params");
assert_eq!(NvptxParamSpace::Registers.to_string(), "registers");
assert_eq!(NvptxParamSpace::Stack.to_string(), "stack");
}
#[test]
fn test_addr_space_ptx_modifiers() {
assert_eq!(NvptxAddrSpace::Global.ptx_modifier(), ".global");
assert_eq!(NvptxAddrSpace::Shared.ptx_modifier(), ".shared");
assert_eq!(NvptxAddrSpace::Const.ptx_modifier(), ".const");
assert_eq!(NvptxAddrSpace::Local.ptx_modifier(), ".local");
assert_eq!(NvptxAddrSpace::Param.ptx_modifier(), ".param");
assert_eq!(NvptxAddrSpace::Reg.ptx_modifier(), ".reg");
}
#[test]
fn test_addr_space_writable() {
assert!(NvptxAddrSpace::Global.is_writable());
assert!(NvptxAddrSpace::Shared.is_writable());
assert!(!NvptxAddrSpace::Const.is_writable());
assert!(!NvptxAddrSpace::Param.is_writable());
}
#[test]
fn test_addr_space_cached() {
assert!(NvptxAddrSpace::Shared.is_cached());
assert!(NvptxAddrSpace::Const.is_cached());
assert!(NvptxAddrSpace::Texture.is_cached());
assert!(!NvptxAddrSpace::Local.is_cached());
}
#[test]
fn test_addr_space_to_llvm() {
assert_eq!(NvptxAddrSpace::Global.to_llvm_addrspace(), 1);
assert_eq!(NvptxAddrSpace::Shared.to_llvm_addrspace(), 3);
assert_eq!(NvptxAddrSpace::Const.to_llvm_addrspace(), 4);
assert_eq!(NvptxAddrSpace::Local.to_llvm_addrspace(), 5);
assert_eq!(NvptxAddrSpace::Param.to_llvm_addrspace(), 101);
}
#[test]
fn test_addr_space_from_llvm() {
assert_eq!(
NvptxAddrSpace::from_llvm_addrspace(1),
Some(NvptxAddrSpace::Global)
);
assert_eq!(
NvptxAddrSpace::from_llvm_addrspace(3),
Some(NvptxAddrSpace::Shared)
);
assert_eq!(
NvptxAddrSpace::from_llvm_addrspace(4),
Some(NvptxAddrSpace::Const)
);
assert_eq!(
NvptxAddrSpace::from_llvm_addrspace(101),
Some(NvptxAddrSpace::Param)
);
assert_eq!(NvptxAddrSpace::from_llvm_addrspace(99), None);
}
#[test]
fn test_instruction_mapping_direct() {
let m = NvptxInstructionMapping::new(NvptxOpcode::Add, vec![NvptxX86Opcode::Add], 1);
assert!(m.is_direct);
assert_eq!(m.expansion_factor, 1);
assert!(m.preserves_semantics);
}
#[test]
fn test_instruction_mapping_expanded() {
let m = NvptxInstructionMapping::new(
NvptxOpcode::Mad,
vec![NvptxX86Opcode::IMul, NvptxX86Opcode::Add],
2,
);
assert!(!m.is_direct);
assert_eq!(m.expansion_factor, 2);
}
#[test]
fn test_instruction_mapping_with_simt() {
let m = NvptxInstructionMapping::new(NvptxOpcode::Bra, vec![NvptxX86Opcode::Jmp], 1)
.with_simt(true);
assert!(m.simt_aware);
}
#[test]
fn test_register_info_x86_new() {
let info = NvptxRegisterInfoX86::new();
assert_eq!(info.allocated_virtual_count(), 0);
assert!(info.mapped_special_count() > 0);
}
#[test]
fn test_register_info_allocate_virtual() {
let mut info = NvptxRegisterInfoX86::new();
let (id, reg) = info.allocate_virtual_reg(64);
assert_eq!(id, 0);
assert_eq!(reg, NvptxX86Reg::RAX);
assert_eq!(info.allocated_virtual_count(), 1);
}
#[test]
fn test_register_info_allocate_multiple() {
let mut info = NvptxRegisterInfoX86::new();
info.allocate_virtual_reg(64);
info.allocate_virtual_reg(64);
info.allocate_virtual_reg(64);
assert_eq!(info.allocated_virtual_count(), 3);
}
#[test]
fn test_register_info_allocate_wide() {
let mut info = NvptxRegisterInfoX86::new();
let (id, reg) = info.allocate_virtual_reg(256);
assert_eq!(id, 0);
assert_eq!(reg, NvptxX86Reg::YMM0);
}
#[test]
fn test_register_info_lookup() {
let mut info = NvptxRegisterInfoX86::new();
let (id, _) = info.allocate_virtual_reg(64);
let looked_up = info.lookup_x86(id);
assert!(looked_up.is_some());
}
#[test]
fn test_register_info_lookup_special() {
let info = NvptxRegisterInfoX86::new();
let reg = info.lookup_special_x86(NvptxSpecialReg::TidX);
assert_eq!(reg, Some(NvptxX86Reg::RAX));
}
#[test]
fn test_register_info_lookup_nonexistent() {
let info = NvptxRegisterInfoX86::new();
assert_eq!(info.lookup_x86(9999), None);
}
#[test]
fn test_register_info_free() {
let mut info = NvptxRegisterInfoX86::new();
let (id, _) = info.allocate_virtual_reg(64);
assert_eq!(info.allocated_virtual_count(), 1);
info.free_virtual_reg(id);
assert_eq!(info.allocated_virtual_count(), 0);
}
#[test]
fn test_register_info_reset() {
let mut info = NvptxRegisterInfoX86::new();
info.allocate_virtual_reg(64);
info.allocate_virtual_reg(128);
assert_eq!(info.allocated_virtual_count(), 2);
info.reset();
assert_eq!(info.allocated_virtual_count(), 0);
}
#[test]
fn test_target_machine_new() {
let tm = NvptxX86TargetMachine::new(75);
assert_eq!(tm.sm_version, 75);
assert_eq!(tm.opt_level, 2);
assert!(tm.mapping_count() > 0);
}
#[test]
fn test_target_machine_opt_level() {
let mut tm = NvptxX86TargetMachine::new(80);
tm.set_opt_level(3);
assert_eq!(tm.opt_level, 3);
tm.set_opt_level(10);
assert_eq!(tm.opt_level, 3);
}
#[test]
fn test_target_machine_avx512() {
let mut tm = NvptxX86TargetMachine::new(80);
assert!(!tm.use_avx512);
tm.enable_avx512();
assert!(tm.use_avx512);
}
#[test]
fn test_target_machine_features() {
let tm = NvptxX86TargetMachine::new(80);
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_features_avx512() {
let mut tm = NvptxX86TargetMachine::new(80);
tm.enable_avx512();
let features = tm.x86_feature_string();
assert!(features.contains("avx512f"));
}
#[test]
fn test_target_machine_lookup_mapping() {
let tm = NvptxX86TargetMachine::new(80);
let m = tm.lookup_mapping(NvptxOpcode::Add);
assert!(m.is_some());
assert_eq!(m.unwrap().expansion_factor, 1);
}
#[test]
fn test_target_machine_lookup_nonexistent() {
let tm = NvptxX86TargetMachine::new(80);
let m = tm.lookup_mapping(NvptxOpcode::Suq); let s = m.map(|x| x.expansion_factor);
assert_eq!(s, Some(1));
}
#[test]
fn test_target_machine_mapping_count_at_least_80() {
let tm = NvptxX86TargetMachine::new(80);
assert!(tm.mapping_count() >= 80);
}
#[test]
fn test_target_machine_compute_capability() {
let tm = NvptxX86TargetMachine::new(75);
assert!((tm.compute_capability() - 7.5).abs() < 0.01);
}
#[test]
fn test_target_machine_data_layout() {
let tm = NvptxX86TargetMachine::new(80);
let dl = tm.get_x86_data_layout();
assert!(dl.starts_with("e-m:e"));
}
#[test]
fn test_bridge_creation() {
let bridge = NVPTXX86Bridge::new(75);
assert!(bridge.is_ready());
assert!(bridge.is_64bit);
assert_eq!(bridge.sm_version(), 75);
assert_eq!(bridge.calling_convention, NvptxCallingConvention::Kernel);
}
#[test]
fn test_bridge_calling_convention() {
let mut bridge = NVPTXX86Bridge::new(80);
bridge.set_calling_convention(NvptxCallingConvention::DeviceFunction);
assert_eq!(
bridge.calling_convention,
NvptxCallingConvention::DeviceFunction
);
}
#[test]
fn test_bridge_32bit() {
let mut bridge = NVPTXX86Bridge::new(75);
bridge.set_64bit(false);
assert!(!bridge.is_64bit);
let triple = bridge.get_nvptx_triple();
assert!(triple.contains("nvptx-"));
assert!(!triple.contains("nvptx64"));
}
#[test]
fn test_bridge_kernel_params() {
let mut bridge = NVPTXX86Bridge::new(75);
bridge.add_kernel_param(NvptxKernelParam::new("a", 4, 4, 0));
bridge.add_kernel_param(NvptxKernelParam::new("b", 8, 8, 4));
assert_eq!(bridge.kernel_params.len(), 2);
bridge.clear_kernel_params();
assert_eq!(bridge.kernel_params.len(), 0);
}
#[test]
fn test_bridge_translate_instruction() {
let bridge = NVPTXX86Bridge::new(80);
let result = bridge.translate_instruction(NvptxOpcode::Add);
assert!(result.is_some());
let (ops, exp) = result.unwrap();
assert_eq!(exp, 1);
assert_eq!(ops[0], NvptxX86Opcode::Add);
}
#[test]
fn test_bridge_direct_mapping() {
let bridge = NVPTXX86Bridge::new(80);
assert!(bridge.has_direct_mapping(NvptxOpcode::Add));
assert!(!bridge.has_direct_mapping(NvptxOpcode::Mad));
}
#[test]
fn test_bridge_compute_expansion() {
let bridge = NVPTXX86Bridge::new(80);
let ops = &[NvptxOpcode::Add, NvptxOpcode::Mad, NvptxOpcode::Sub];
let exp = bridge.compute_expansion(ops);
assert_eq!(exp, 4); }
#[test]
fn test_bridge_avx512() {
let mut bridge = NVPTXX86Bridge::new(80);
bridge.enable_avx512();
assert!(bridge.stats.avx512_enabled);
}
#[test]
fn test_bridge_opt_level() {
let mut bridge = NVPTXX86Bridge::new(80);
bridge.set_opt_level(3);
assert_eq!(bridge.target_machine.opt_level, 3);
}
#[test]
fn test_bridge_x86_prologue() {
let bridge = NVPTXX86Bridge::new(75);
let p = bridge.emit_x86_prologue();
assert!(p.contains("NVPTX Kernel Prologue"));
assert!(p.contains("SM Version: sm_75"));
assert!(p.contains("Compute Capability: 7.5"));
}
#[test]
fn test_bridge_x86_epilogue() {
let bridge = NVPTXX86Bridge::new(75);
let ep = bridge.emit_x86_epilogue();
assert!(ep.contains("NVPTX Kernel Epilogue"));
}
#[test]
fn test_bridge_x86_prologue_with_params() {
let mut bridge = NVPTXX86Bridge::new(80);
bridge.add_kernel_param(NvptxKernelParam::new("input", 4, 4, 0));
let p = bridge.emit_x86_prologue();
assert!(p.contains("input"));
assert!(p.contains("Kernel Parameters"));
}
#[test]
fn test_bridge_nvptx_triple() {
let bridge = NVPTXX86Bridge::new(80);
assert_eq!(bridge.get_nvptx_triple(), "nvptx64-nvidia-cuda");
}
#[test]
fn test_bridge_nvptx_triple_32() {
let mut bridge = NVPTXX86Bridge::new(80);
bridge.set_64bit(false);
assert_eq!(bridge.get_nvptx_triple(), "nvptx-nvidia-cuda");
}
#[test]
fn test_bridge_allocate_virtual_reg() {
let mut bridge = NVPTXX86Bridge::new(80);
let (id, reg) = bridge.allocate_virtual_reg(64);
assert_eq!(id, 0);
assert_eq!(reg, NvptxX86Reg::RAX);
}
#[test]
fn test_bridge_stats() {
let bridge = NVPTXX86Bridge::new(80);
let s = bridge.get_stats();
assert_eq!(s.instructions_translated, 0);
assert_eq!(s.kernels_processed, 0);
assert!(!s.avx512_enabled);
}
#[test]
fn test_bridge_stats_avx512() {
let mut bridge = NVPTXX86Bridge::new(80);
bridge.enable_avx512();
assert!(bridge.stats.avx512_enabled);
}
#[test]
fn test_bridge_reset_stats() {
let mut bridge = NVPTXX86Bridge::new(80);
bridge.enable_avx512();
bridge.reset_stats();
assert!(!bridge.stats.avx512_enabled);
}
#[test]
fn test_cross_target_new() {
let ct = CrossTargetNvptx::new(80);
assert_eq!(ct.warp_size, 32);
assert_eq!(ct.max_threads_per_block, 1024);
assert!(ct.scalarize_uniform);
assert!(ct.emit_lane_masks);
}
#[test]
fn test_cross_target_lane_regs() {
let mut ct = CrossTargetNvptx::new(80);
ct.set_exec_mask_reg(NvptxX86Reg::RAX);
assert_eq!(ct.exec_mask_reg, Some(NvptxX86Reg::RAX));
ct.set_lane_id_reg(NvptxX86Reg::RBX);
assert_eq!(ct.lane_id_reg, Some(NvptxX86Reg::RBX));
}
#[test]
fn test_cross_target_full_warp_mask() {
assert_eq!(CrossTargetNvptx::full_warp_mask(), 0xFFFF_FFFF);
}
#[test]
fn test_cross_target_warps_per_block() {
let ct = CrossTargetNvptx::new(80);
assert_eq!(ct.warps_per_block(32), 1);
assert_eq!(ct.warps_per_block(64), 2);
assert_eq!(ct.warps_per_block(33), 2);
assert_eq!(ct.warps_per_block(1024), 32);
}
#[test]
fn test_cross_target_threads_from_warps() {
let ct = CrossTargetNvptx::new(80);
assert_eq!(ct.threads_from_warps(1), 32);
assert_eq!(ct.threads_from_warps(4), 128);
}
#[test]
fn test_cross_target_emit_lane_init() {
let ct = CrossTargetNvptx::new(80);
let code = ct.emit_lane_init();
assert!(code.contains("SIMT Lane Initialization"));
assert!(code.contains("32"));
assert!(code.contains("1024"));
}
#[test]
fn test_cross_target_emit_warp_loop() {
let ct = CrossTargetNvptx::new(80);
let code = ct.emit_warp_loop_header();
assert!(code.contains("Warp Loop Header"));
}
#[test]
fn test_cross_target_uniform_instruction() {
assert!(CrossTargetNvptx::is_uniform_instruction(NvptxOpcode::Bar));
assert!(CrossTargetNvptx::is_uniform_instruction(
NvptxOpcode::BarSync
));
assert!(CrossTargetNvptx::is_uniform_instruction(
NvptxOpcode::Membar
));
assert!(!CrossTargetNvptx::is_uniform_instruction(NvptxOpcode::Add));
assert!(!CrossTargetNvptx::is_uniform_instruction(NvptxOpcode::Mov));
}
#[test]
fn test_cross_target_reset() {
let mut ct = CrossTargetNvptx::new(80);
ct.set_lane_id_reg(NvptxX86Reg::RAX);
ct.set_exec_mask_reg(NvptxX86Reg::RBX);
ct.active_warps = 4;
ct.reset();
assert_eq!(ct.lane_id_reg, None);
assert_eq!(ct.exec_mask_reg, None);
assert_eq!(ct.active_warps, 0);
}
#[test]
fn test_kernel_descriptor_new() {
let kd = NvptxKernelDescriptorX86::new("test");
assert_eq!(kd.name, "test");
assert_eq!(kd.register_count, 0);
assert_eq!(kd.shared_memory, 0);
assert!(!kd.uses_dynamic_shared);
assert_eq!(kd.max_threads_per_block, 1024);
}
#[test]
fn test_kernel_descriptor_with_params() {
let mut kd = NvptxKernelDescriptorX86::new("k");
kd.add_param(NvptxKernelParam::new("a", 4, 4, 0));
kd.add_param(NvptxKernelParam::new("b", 8, 8, 4));
assert_eq!(kd.params.len(), 2);
}
#[test]
fn test_kernel_descriptor_register_usage() {
let mut kd = NvptxKernelDescriptorX86::new("k");
kd.set_register_usage(32);
assert_eq!(kd.register_count, 32);
}
#[test]
fn test_kernel_descriptor_shared_memory() {
let mut kd = NvptxKernelDescriptorX86::new("k");
kd.set_shared_memory(4096, true);
assert_eq!(kd.shared_memory, 4096);
assert!(kd.uses_dynamic_shared);
}
#[test]
fn test_kernel_descriptor_emit_launch_stub() {
let mut kd = NvptxKernelDescriptorX86::new("my_kernel");
kd.set_register_usage(16);
kd.set_shared_memory(2048, false);
let stub = kd.emit_x86_launch_stub();
assert!(stub.contains("my_kernel"));
assert!(stub.contains("X86 Launch Stub"));
assert!(stub.contains("shared_mem=2048"));
}
#[test]
fn test_kernel_descriptor_emit_ptx_entry() {
let mut kd = NvptxKernelDescriptorX86::new("my_kernel");
kd.add_param(NvptxKernelParam::new("a", 4, 4, 0));
kd.add_param(NvptxKernelParam::new("b", 8, 8, 4));
kd.set_register_usage(24);
let ptx = kd.emit_ptx_entry();
assert!(ptx.contains(".visible .entry my_kernel"));
assert!(ptx.contains(".param"));
assert!(ptx.contains(".maxnreg 24"));
}
#[test]
fn test_kernel_param_new() {
let p = NvptxKernelParam::new("a", 4, 4, 0);
assert_eq!(p.name, "a");
assert_eq!(p.size_bytes, 4);
assert!(!p.is_pointer);
assert_eq!(p.addr_space, NvptxAddrSpace::Param);
}
#[test]
fn test_kernel_param_with_pointer() {
let p = NvptxKernelParam::new("p", 8, 8, 0).with_pointer(true);
assert!(p.is_pointer);
}
#[test]
fn test_kernel_param_with_addr_space() {
let p = NvptxKernelParam::new("s", 4, 4, 0).with_addr_space(NvptxAddrSpace::Shared);
assert_eq!(p.addr_space, NvptxAddrSpace::Shared);
}
#[test]
fn test_memory_order_ptx_suffix() {
assert_eq!(NvptxMemoryOrder::Relaxed.to_ptx_suffix(), ".relaxed");
assert_eq!(NvptxMemoryOrder::Acquire.to_ptx_suffix(), ".acquire");
assert_eq!(NvptxMemoryOrder::Release.to_ptx_suffix(), ".release");
assert_eq!(NvptxMemoryOrder::AcquireRelease.to_ptx_suffix(), ".acq_rel");
assert_eq!(
NvptxMemoryOrder::SequentiallyConsistent.to_ptx_suffix(),
".sc"
);
}
#[test]
fn test_memory_order_x86_barrier() {
assert_eq!(NvptxMemoryOrder::Relaxed.to_x86_barrier(), None);
assert_eq!(NvptxMemoryOrder::Acquire.to_x86_barrier(), None);
assert_eq!(
NvptxMemoryOrder::SequentiallyConsistent.to_x86_barrier(),
Some("mfence")
);
}
#[test]
fn test_memory_scope_ptx_suffix() {
assert_eq!(NvptxMemoryScope::Cta.to_ptx_suffix(), ".cta");
assert_eq!(NvptxMemoryScope::Gpu.to_ptx_suffix(), ".gpu");
assert_eq!(NvptxMemoryScope::Sys.to_ptx_suffix(), ".sys");
}
#[test]
fn test_memory_scope_x86_barrier() {
assert_eq!(NvptxMemoryScope::Cta.to_x86_barrier(), None);
assert_eq!(NvptxMemoryScope::Gpu.to_x86_barrier(), Some("mfence"));
}
#[test]
fn test_warp_emulator_new() {
let we = NvptxWarpEmulator::new();
assert_eq!(we.warp_size, 32);
assert!(we.all_active());
}
#[test]
fn test_warp_emulator_lane_mask() {
let mut we = NvptxWarpEmulator::new();
we.deactivate_lane(0);
assert!(!we.all_active());
we.activate_lane(0);
assert!(we.all_active());
}
#[test]
fn test_warp_emulator_emit_shfl() {
let we = NvptxWarpEmulator::new();
let code = we.emit_shfl_sync(0, 32);
assert!(code.contains("shfl.sync"));
assert!(code.contains("lane 0"));
}
#[test]
fn test_builder_default() {
let builder = NvptxBridgeBuilder::new(80);
let bridge = builder.build();
assert!(bridge.is_ready());
assert!(bridge.is_64bit);
}
#[test]
fn test_builder_32bit() {
let builder = NvptxBridgeBuilder::new(80).bit32();
let bridge = builder.build();
assert!(!bridge.is_64bit);
}
#[test]
fn test_builder_with_options() {
let builder = NvptxBridgeBuilder::new(75)
.opt_level(3)
.avx512(true)
.simt_emulation(false)
.calling_convention(NvptxCallingConvention::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,
NvptxCallingConvention::DeviceFunction
);
}
#[test]
fn test_builder_with_params() {
let builder =
NvptxBridgeBuilder::new(80).add_kernel_param(NvptxKernelParam::new("a", 4, 4, 0));
let bridge = builder.build();
assert_eq!(bridge.kernel_params.len(), 1);
}
#[test]
fn test_builder_avx_flags() {
let bridge = NvptxBridgeBuilder::new(80).avx(false).avx2(false).build();
assert!(!bridge.target_machine.use_avx);
assert!(!bridge.target_machine.use_avx2);
}
#[test]
fn test_all_sm_versions_build() {
for sm in &[50u32, 52, 53, 60, 61, 62, 70, 72, 75, 80, 86, 87, 89, 90] {
let bridge = NVPTXX86Bridge::new(*sm);
assert!(bridge.is_ready(), "Failed for SM {}", sm);
}
}
#[test]
fn test_instruction_mapping_count_at_least_80() {
let tm = NvptxX86TargetMachine::new(80);
assert!(tm.mapping_count() >= 80);
}
#[test]
fn test_register_info_all_special_mapped() {
let info = NvptxRegisterInfoX86::new();
assert_eq!(info.mapped_special_count(), 10);
}
#[test]
fn test_addr_space_display() {
assert_eq!(NvptxAddrSpace::Global.to_string(), ".global");
assert_eq!(NvptxAddrSpace::Shared.to_string(), ".shared");
}
#[test]
fn test_reg_class_display() {
assert_eq!(NvptxX86RegClass::GPR.to_string(), "GPR");
assert_eq!(NvptxX86RegClass::XMM.to_string(), "XMM");
}
#[test]
fn test_special_reg_unmapped_runtime() {
assert_eq!(NvptxSpecialReg::SmId.map_to_x86_runtime(), None);
assert_eq!(NvptxSpecialReg::Clock.map_to_x86_runtime(), None);
}
#[test]
fn test_bridge_implicit_convention_params() {
assert_eq!(NvptxCallingConvention::Kernel.implicit_param_count(), 0);
assert_eq!(
NvptxCallingConvention::DeviceFunction.implicit_param_count(),
0
);
}
#[test]
fn test_warp_emulator_ballot() {
let we = NvptxWarpEmulator::new();
let code = we.emit_ballot();
assert!(code.contains("ballot"));
}
#[test]
fn test_cross_target_exec_predicate_inv() {
let ct = CrossTargetNvptx::new(80);
let code = ct.emit_exec_predicate(true);
assert!(code.contains("skip"));
}
#[test]
fn test_kernel_descriptor_min_threads() {
let mut kd = NvptxKernelDescriptorX86::new("k");
kd.min_threads_per_block = 64;
assert_eq!(kd.min_threads_per_block, 64);
}
}