pub enum MatmulPrecision {
Auto,
Int8Dp4a,
Int64,
Bf16Hmma,
Fp16Hmma,
}Expand description
Contraction precision for a lowered matmul.
Variants§
Auto
Best general case: pick the always-correct path unless the value analysis proves a faster one is
safe. Today resolves to ResolvedPrecision::Int64 (we do not yet carry a proof that operands fit
int8 through the recognizer), so Auto is exactly “the correct-for-any-program default”.
Int8Dp4a
Signed-int8 __dp4a accumulation — faster, but VALID ONLY when operands provably fit int8 and the
contraction fits int32. Not strictly better (restricted domain) → a knob, off by default.
Int64
64-bit integer accumulation — always correct (matches the VM’s i64 semantics for any operands), slower.
Bf16Hmma
bf16 tensor-core GEMM (HMMA 16×16×16, bf16 operands, f32 accumulate) — the half-precision fast path
for a matmul over Real tensors on a 16-aligned shape (sm_80+). When the block tile divides the shape
and its footprint fits static shared memory, the lowering auto-selects the cp.async double-buffered
kernel (the measured 1.43× lever, select_hmma_kernel); otherwise the basic one-warp WMMA kernel. On an
incapable shape/arch it falls back to the always-correct int64 kernel. A knob (restricted domain), off by default.
Fp16Hmma
fp16 tensor-core GEMM (HMMA 16×16×16, f16 operands, f32 accumulate) — the fp16 counterpart (sm_70+).
Implementations§
Source§impl MatmulPrecision
impl MatmulPrecision
Sourcepub fn resolved(self) -> ResolvedPrecision
pub fn resolved(self) -> ResolvedPrecision
Resolve the knob to a concrete lowering. Auto conservatively picks the always-correct Int64 — we
only emit a fast path (Int8Dp4a/Bf16Hmma/Fp16Hmma) when the caller explicitly requests it,
which is precisely “default to best general; opt into the trade-off”.
Trait Implementations§
Source§impl Clone for MatmulPrecision
impl Clone for MatmulPrecision
Source§fn clone(&self) -> MatmulPrecision
fn clone(&self) -> MatmulPrecision
1.0.0 · Source§fn clone_from(&mut self, source: &Self)
fn clone_from(&mut self, source: &Self)
source. Read moreSource§impl Debug for MatmulPrecision
impl Debug for MatmulPrecision
Source§impl PartialEq for MatmulPrecision
impl PartialEq for MatmulPrecision
impl Copy for MatmulPrecision
impl Eq for MatmulPrecision
impl StructuralPartialEq for MatmulPrecision
Auto Trait Implementations§
impl Freeze for MatmulPrecision
impl RefUnwindSafe for MatmulPrecision
impl Send for MatmulPrecision
impl Sync for MatmulPrecision
impl Unpin for MatmulPrecision
impl UnsafeUnpin for MatmulPrecision
impl UnwindSafe for MatmulPrecision
Blanket Implementations§
Source§impl<T> BorrowMut<T> for Twhere
T: ?Sized,
impl<T> BorrowMut<T> for Twhere
T: ?Sized,
Source§fn borrow_mut(&mut self) -> &mut T
fn borrow_mut(&mut self) -> &mut T
Source§impl<T> CloneToUninit for Twhere
T: Clone,
impl<T> CloneToUninit for Twhere
T: Clone,
§impl<T> Downcast for Twhere
T: Any,
impl<T> Downcast for Twhere
T: Any,
§fn into_any(self: Box<T>) -> Box<dyn Any>
fn into_any(self: Box<T>) -> Box<dyn Any>
Box<dyn Trait> (where Trait: Downcast) to Box<dyn Any>. Box<dyn Any> can
then be further downcast into Box<ConcreteType> where ConcreteType implements Trait.§fn into_any_rc(self: Rc<T>) -> Rc<dyn Any>
fn into_any_rc(self: Rc<T>) -> Rc<dyn Any>
Rc<Trait> (where Trait: Downcast) to Rc<Any>. Rc<Any> can then be
further downcast into Rc<ConcreteType> where ConcreteType implements Trait.§fn as_any(&self) -> &(dyn Any + 'static)
fn as_any(&self) -> &(dyn Any + 'static)
&Trait (where Trait: Downcast) to &Any. This is needed since Rust cannot
generate &Any’s vtable from &Trait’s.§fn as_any_mut(&mut self) -> &mut (dyn Any + 'static)
fn as_any_mut(&mut self) -> &mut (dyn Any + 'static)
&mut Trait (where Trait: Downcast) to &Any. This is needed since Rust cannot
generate &mut Any’s vtable from &mut Trait’s.§impl<T> DowncastSync for T
impl<T> DowncastSync for T
§impl<Q, K> Equivalent<K> for Q
impl<Q, K> Equivalent<K> for Q
§fn equivalent(&self, key: &K) -> bool
fn equivalent(&self, key: &K) -> bool
§impl<Q, K> Equivalent<K> for Q
impl<Q, K> Equivalent<K> for Q
§fn equivalent(&self, key: &K) -> bool
fn equivalent(&self, key: &K) -> bool
key and return true if they are equal.