Enum RejectReason 

pub enum RejectReason {
    Atomic(&'static str),
    StrictFp(&'static str),
    TableOp(&'static str),
    GcOp(&'static str),
    MemoryResize(&'static str),
    LargeMemcpy(&'static str),
    HostCall(&'static str),
    FloatOp(&'static str),
    IntDivRem(&'static str),
    BackEdge(&'static str),
    Trap(&'static str),
}

A reason a Cranelift instruction cannot be lowered to a crate::lowered_ir::LoweredOp.

Each variant carries the offending opcode’s mnemonic (&'static str) so error messages and logs are grep-able against the mapping table. The mnemonic is the same string Cranelift’s own Opcode Display impl prints (e.g. "atomic_rmw", not "AtomicRmw"), making it stable across the cranelift-codegen version bumps that periodically rename enum variants but keep the textual mnemonic.

Variants

Atomic(&'static str)

Atomic memory operation. Deferred — Wasm-threads-on-GPU is a memory-model problem out of scope for wave 2. Covers atomic_load, atomic_store, atomic_rmw, atomic_cas.

StrictFp(&'static str)

Floating-point opcode with strict-FP exception semantics PTX default rounding cannot match. Covers fcvt_to_sint_sat and fcvt_to_uint_sat — both saturate-on-out-of-range, behaviour the PTX cvt does not provide by default. The full strict-FP scope (NaN propagation, denormals-as-zero) is broader; wave 2 limits itself to the trap-carrying conversions.

TableOp(&'static str)

Wasm table.get / table.set. Tables live host-side; lowering them would require device-resident table mirrors. Hard-rejected.

Not currently wired in the detector: cranelift-codegen 0.111 has no direct Opcode::TableGet / Opcode::TableSet variant — cranelift-wasm lowers these to load/store+libcall sequences in the translator, so a Wasm module reaches us with the table access already expanded into a call to a runtime helper (and therefore is caught by RejectReason::HostCall). Pre-declared here so a future direct Wasm front-end or a Cranelift bump that introduces table opcodes natively does not have to reshuffle the public enum.

GcOp(&'static str)

Wasm GC / reference-type opcode (ref.func, ref.null, ref.is_null on a non-i31 ref). No device-side representation possible. Hard-rejected.

Not currently wired for the same reason as RejectReason::TableOp: in cranelift-codegen 0.111 the GC proposal opcodes either do not exist as direct enum variants (ref.func) or are translated into other ops by cranelift-wasm before the detector sees them. Pre-declared for forward compatibility.

MemoryResize(&'static str)

memory.grow / memory.size. Linear-memory resizing requires a host round-trip; PTX kernels must run with a fixed memory snapshot.

Not currently wired: cranelift-codegen 0.111 has no direct Opcode::MemoryGrow / Opcode::MemorySize — Wasm memory.grow reaches us as a libcall (a Call instruction targeting the runtime helper) and is therefore caught by RejectReason::HostCall. Pre-declared for the same forward- compatibility reason as the other absent categories.

LargeMemcpy(&'static str)

memory.copy / memory.fill above the inline-copy threshold. PTX has cp.async.bulk (sm_90+) but the wave-2 baseline is sm_80 — the cuda-oxide PTX target version pinned via crate::ptx_emit::DEFAULT_TARGET.

Not currently wired: same translator-expansion situation as the other absent categories — memory.copy / memory.fill reach us as Call instructions to the runtime, caught by RejectReason::HostCall. Pre-declared for forward compatibility. Wave 3+ will refine RejectReason::HostCall to distinguish device-internal calls from runtime-helper calls and may at that point route small memory.copy libcalls back to an inline Load/Store sequence instead of a hard reject.

HostCall(&'static str)

call / call_indirect / return_call / return_call_indirect. Host-callback prohibition: PTX has no path back into the Wasmtime runtime, so every call is rejected at the wave-2 detector. Wave 3+ will distinguish device-internal calls (legal: lowered to func.call) from host round-trips (illegal: rejected).

FloatOp(&'static str)

Any floating-point opcode (fadd, fmul, fdiv, fma, sqrt, fcmp, the float conversions, …).

jit MED fix (finding 4): the reject-list previously admitted non-saturating FP ops. PTX default rounding (add.rn.f32 etc.) is bit-exact IEEE for the basic ops, but the broader strict-FP scope (NaN payload propagation, denormal-as-zero behaviour, transcendental approximations) is NOT yet proven equivalent to the Wasm/CPU reference across this lowering path. Until a differential proof of bit-exact rounding lands, ALL float opcodes are rejected so a float function deopts to the CPU path rather than risk silent numerical divergence. (The basic-op subset can be re-admitted once the differential oracle proves equivalence — narrow this then.)

IntDivRem(&'static str)

Integer divide / remainder (sdiv, udiv, srem, urem).

jit MED fix (finding 4): Wasm i32.div_u etc. TRAP on divide-by- zero (and signed div traps on INT_MIN / -1 overflow). PTX div has undefined behaviour on divide-by-zero — it does not trap. Until the lowering emits an explicit divisor-zero guard that reproduces the Wasm trap, integer div/rem is rejected so it stays on the CPU path rather than producing a non-trapping (wrong) result on the GPU.

BackEdge(&'static str)

A control-flow back-edge: a branch whose target is a block at or before the branching block in layout order (i.e. a loop).

jit MED fix (finding 4): the wave-2 lowering does not yet model loop carried dependencies / convergence on the GPU, so any function containing a loop back-edge is rejected. Straight-line and forward-only (if/else, switch) control flow remains admissible.

Trap(&'static str)

An explicit trap or unreachable (trap, trapz, trapnz, resumable_trap, debugtrap; Wasm unreachable lowers to trap).

jit MED fix (finding 4): PTX has no equivalent of the Wasm trap machinery (which unwinds back into the host with a trap code). A function that can trap mid-kernel cannot be faithfully offloaded, so it is rejected.

Implementations

impl RejectReason

pub fn opcode_mnemonic(&self) -> &'static str

The Cranelift opcode mnemonic that triggered this rejection.

Convenience accessor for log lines / error formatting that want the offending opcode name without match-ing on the variant.

Trait Implementations

impl Clone for RejectReason

fn clone(&self) -> RejectReason

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for RejectReason

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Eq for RejectReason

impl PartialEq for RejectReason

fn eq(&self, other: &RejectReason) -> bool

Tests for self and other values to be equal, and is used by ==.

fn ne(&self, other: &Rhs) -> bool

Tests for !=. The default implementation is almost always sufficient, and should not be overridden without very good reason.

impl StructuralPartialEq for RejectReason