Struct Fp8E4M3 

#[repr(transparent)]
pub struct Fp8E4M3(pub u8);

8-bit floating-point, E4M3 encoding (1 sign + 4 exponent + 3 mantissa, exponent bias 7).

#[repr(transparent)] around u8 storage — bit-compatible with __nv_fp8_storage_t on the CUDA side and with float8::F8E4M3 on the host side. A DeviceBuffer<u8> (byte substrate) can be reinterpreted as DeviceBuffer<Fp8E4M3> via view_as without copying.

Numerical range: ±448 (max finite). One NaN encoding only (S.1111.111); E4M3 has no infinities. The conversion path matches NVIDIA’s __nv_cvt_float_to_fp8(x, __NV_SATFINITE, __NV_E4M3): round-half-to-even, saturating-to-max-finite on overflow.

Routes through Ada Lovelace FP8 tensor cores (mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32) with F32 accumulation and float alpha / beta scaling. First landed in baracuda-kernels Phase 2.

Tuple Fields

0: u8

Implementations

impl Fp8E4M3

pub fn from_f32(x: f32) -> Fp8E4M3

Convert from f32 using NVIDIA’s SATFINITE semantics (round-half-to-even, clamp |x| to the E4M3 max-finite 448.0).

pub fn to_f32(self) -> f32

Convert to f32. The E4M3 grid is sparse — the result is one of 254 finite values (or NaN) on the E4M3 lattice.

Trait Implementations

impl Clone for Fp8E4M3

fn clone(&self) -> Fp8E4M3

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Copy for Fp8E4M3

impl Debug for Fp8E4M3

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

Formats the value using the given formatter.

impl Default for Fp8E4M3

fn default() -> Fp8E4M3

Returns the “default value” for a type.

impl DeviceRepr for Fp8E4M3

impl Eq for Fp8E4M3

impl FpElement for Fp8E4M3

impl Hash for Fp8E4M3

fn hash<__H>(&self, state: &mut __H)

where __H: Hasher,

Feeds this value into the given Hasher.

fn hash_slice<H>(data: &[Self], state: &mut H)

where H: Hasher, Self: Sized,

Feeds a slice of this type into the given Hasher.

impl KernelDtype for Fp8E4M3

const KIND: ElementKind = ElementKind::Fp8E4M3

Runtime tag for this dtype. Stable across the workspace — keyed by this same enum in crate::KernelSku::element.

impl Ord for Fp8E4M3

fn cmp(&self, other: &Fp8E4M3) -> Ordering

This method returns an Ordering between self and other.

fn max(self, other: Self) -> Self

where Self: Sized,

Compares and returns the maximum of two values.

fn min(self, other: Self) -> Self

where Self: Sized,

Compares and returns the minimum of two values.

fn clamp(self, min: Self, max: Self) -> Self

where Self: Sized,

Restrict a value to a certain interval.

impl PartialEq for Fp8E4M3

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

fn partial_cmp(&self, other: &Fp8E4M3) -> Option<Ordering>

This method returns an ordering between self and other values if one exists.

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

Tests less than (for self and other) and is used by the < operator.

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

Tests less than or equal to (for self and other) and is used by the <= operator.

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

Tests greater than (for self and other) and is used by the > operator.

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

Tests greater than or equal to (for self and other) and is used by the >= operator.

impl StructuralPartialEq for Fp8E4M3