Struct S4 

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

Signed 4-bit integer element marker — packed-pair storage.

#[repr(transparent)] around u8. One S4 storage slot is one byte and holds two packed s4 elements: the low nibble is the element at even logical index, the high nibble is the element at odd logical index (along the K axis for A/B operands, along the N axis for D output). Sign-extended to s32 on the GPU side via ((s8)(nibble << 4)) >> 4.

A DeviceBuffer<u8> of (M*K)/2 bytes can be reinterpreted as a DeviceBuffer<S4> of (M*K)/2 storage slots via view_as without copying — S4 is byte-storage at the buffer layer, and element count lives at the plan-layer descriptor (M / N / K).

Numerical range per element: [-8, +7]. The plan layer (Int4GemmPlan in baracuda-kernels) takes M, N, K in element counts and leading dimensions in storage-slot (= byte) counts — MatrixRef<S4>::ld therefore equals K / 2 for row-major A with no padding. K must be even (packing is byte- aligned). Routes through Ada Lovelace int4 tensor cores (mma.sync.aligned.m16n8k64.row.col.satfinite.s32.s4.s4.s32) with S32 accumulation and float alpha / beta scaling. First landed in baracuda-kernels Phase 2.

Tuple Fields

0: u8

Implementations

impl S4

pub fn pack(lo: i8, hi: i8) -> S4

Pack two s4 values [lo, hi] (each in [-8, +7]) into one storage slot. Values outside the range are masked to their low 4 bits (no saturation — pre-clamp on the caller side if needed).

pub fn unpack(self) -> [i8; 2]

Unpack into [low_nibble_as_s4, high_nibble_as_s4]. Each returned value is sign-extended from the 4-bit nibble.

Trait Implementations

impl Clone for S4

fn clone(&self) -> S4

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Copy for S4

impl Debug for S4

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

Formats the value using the given formatter.

impl Default for S4

fn default() -> S4

Returns the “default value” for a type.

impl DeviceRepr for S4

impl Eq for S4

impl Hash for S4

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 IntElement for S4

impl KernelDtype for S4

const KIND: ElementKind = ElementKind::S4

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

impl Ord for S4

fn cmp(&self, other: &S4) -> 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 S4

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

fn partial_cmp(&self, other: &S4) -> 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 S4