Struct Bin 

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

1-bit binary element marker — packed-byte storage.

#[repr(transparent)] around u8. One Bin storage slot is one byte and holds eight packed b1 elements: bit i of the byte (LSB = bit 0) is the element at K offset 8 * byte_idx + i. Packing is along the K axis for A/B operands.

A DeviceBuffer<u8> of (M*K)/8 bytes can be reinterpreted as a DeviceBuffer<Bin> of (M*K)/8 storage slots via view_as without copying.

Routes through Ampere+ binary tensor cores (mma.sync.aligned.m16n8k256.row.col.s32.b1.b1.s32.xor.popc) with an S32 output accumulator. Unlike the int4 / int8 / FP8 families, bin GEMM does not quantize its output back to the input element type — the result is the raw popcount accumulator (popcount(xor(A_row, B_col)) summed over K bytes), surfaced as i32. No α / β / bias / activation chain (the popcount programming model doesn’t have a meaningful place for them).

The plan layer ([Bin is consumed by BinGemmPlan in baracuda-kernels) takes M, N, K in element counts and leading dimensions in storage-slot (= byte) counts — MatrixRef<Bin>::ld therefore equals K / 8 for row-major A with no padding. K must be divisible by 8 (packing is byte-aligned).

Tuple Fields

0: u8

Implementations

impl Bin

pub fn pack(bits: [bool; 8]) -> Bin

Pack eight bool values bits[0..8] into one storage byte. bits[i] becomes bit i of the result (LSB-first).

pub fn unpack(self) -> [bool; 8]

Unpack one storage byte into eight bool values along K (LSB-first).

Trait Implementations

impl BinElement for Bin

impl Clone for Bin

fn clone(&self) -> Bin

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Copy for Bin

impl Debug for Bin

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

Formats the value using the given formatter.

impl Default for Bin

fn default() -> Bin

Returns the “default value” for a type.

impl DeviceRepr for Bin

impl Eq for Bin

impl Hash for Bin

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 Bin

const KIND: ElementKind = ElementKind::Bin

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

impl Ord for Bin

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

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

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