Enum LoweredArg 

pub enum LoweredArg {
    I32(i32),
    I64(i64),
    F32(f32),
    F64(f64),
    U32(u32),
    U64(u64),
    #[non_exhaustive]
    Ptr {
        host_ptr: *const u8,
        len: u32,
        guest_offset: u32,
    },
}

A single fully-parsed and bounds-checked kernel argument.

LoweredArg is the host-side representation that flows into cuLaunchKernel (via build_kernel_param_storage) on CUDA builds and is recorded for inspection on no-CUDA builds.

Variants

I32(i32)

32-bit signed scalar.

I64(i64)

64-bit signed scalar.

F32(f32)

32-bit float scalar.

F64(f64)

64-bit float scalar.

U32(u32)

32-bit unsigned scalar.

U64(u64)

64-bit unsigned scalar.

#[non_exhaustive]

Ptr

Pointer argument: a fully bounds-checked host pointer plus the declared byte length and the original guest-memory offset.

host_ptr is a raw pointer into the caller’s linear memory; under CUDA Unified Memory it is also a valid device address. The guest_offset field is preserved for log lines / tests so a pointer arg round-trips losslessly through parsing.

#[non_exhaustive] blocks external struct-literal construction — out-of-crate callers cannot stamp a LoweredArg::Ptr { ... } expression with an arbitrary host_ptr. The launch path itself, which lives inside this crate, retains full construction rights; integration tests author pointer args via LoweredArg::ptr_for_encoding (null placeholder pointer the parser overwrites with the resolved one). This is defence-in- depth on top of the comment that the field is “treat as opaque”: a guest-driven memory.grow between launch and any embedder readback can dangle the host pointer, so the only safe consumer is the launch path. Embedders observing the parsed argv take a pointer-free LoweredArgSnapshot via crate::host::WasiCudaContext::last_lowered_args.

Pattern-matching still works for external consumers using the .. rest pattern (LoweredArg::Ptr { guest_offset, len, .. }); the rest pattern intentionally hides host_ptr at the public boundary.

Fields

Non-exhaustive enum variants could have additional fields added in future. Therefore, non-exhaustive enum variants cannot be constructed in external crates and cannot be matched against.

host_ptr: *const u8

Host-side raw pointer into the caller’s linear memory.

Treat as opaque — dereferencing is the caller’s responsibility and is unsafe in general (the underlying linear memory may move on a memory.grow). The CUDA launch site captures the pointer into a parameter slot and immediately hands it to cuLaunchKernel, before any guest code can run.

len: u32

Byte length the kernel will access.

guest_offset: u32

Original guest-memory offset (for logging / tests).

Implementations

impl LoweredArg

pub fn ptr_for_encoding(guest_offset: u32, len: u32) -> Self

Test/encoding helper: build a LoweredArg::Ptr carrying a null host pointer.

Out-of-crate callers (integration tests and embedders that want to round-trip argv through encode_argv) cannot construct LoweredArg::Ptr via struct-literal syntax because the variant is #[non_exhaustive]; this constructor is the supported entry point. encode_argv reads only guest_offset and len, so a null placeholder is sufficient for the wire-format encoding path; the launch path itself populates host_ptr via parse_argv.

Trait Implementations

impl Clone for LoweredArg

fn clone(&self) -> LoweredArg

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Debug for LoweredArg

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

Formats the value using the given formatter.

impl From<&LoweredArg> for LoweredArgSnapshot

fn from(arg: &LoweredArg) -> Self

Converts to this type from the input type.

impl PartialEq for LoweredArg

fn eq(&self, other: &LoweredArg) -> 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 Send for LoweredArg

impl StructuralPartialEq for LoweredArg