pub enum HmmaKernel {
CpAsync {
tile: WmmaTile,
stages: usize,
},
Basic,
}Expand description
Which half-precision WMMA GEMM kernel is the FASTEST legal choice for a shape.
Variants§
CpAsync
cp.async double-buffered — the measured 1.43× lever; needs sm_80+, the tile to divide the shape evenly, and the (operand-byte-sized) double-buffered footprint to fit static shared memory.
Basic
The basic one-warp-per-16×16-tile kernel — any 16-aligned shape (the fallback when cp.async is illegal).
Trait Implementations§
Source§impl Clone for HmmaKernel
impl Clone for HmmaKernel
Source§fn clone(&self) -> HmmaKernel
fn clone(&self) -> HmmaKernel
Returns a duplicate of the value. Read more
1.0.0 · Source§fn clone_from(&mut self, source: &Self)
fn clone_from(&mut self, source: &Self)
Performs copy-assignment from
source. Read moreSource§impl Debug for HmmaKernel
impl Debug for HmmaKernel
Source§impl PartialEq for HmmaKernel
impl PartialEq for HmmaKernel
impl Copy for HmmaKernel
impl Eq for HmmaKernel
impl StructuralPartialEq for HmmaKernel
Auto Trait Implementations§
impl Freeze for HmmaKernel
impl RefUnwindSafe for HmmaKernel
impl Send for HmmaKernel
impl Sync for HmmaKernel
impl Unpin for HmmaKernel
impl UnsafeUnpin for HmmaKernel
impl UnwindSafe for HmmaKernel
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
Mutably borrows from an owned value. Read more
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>
Convert
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>
Convert
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)
Convert
&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)
Convert
&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
Checks if this value is equivalent to the given key. Read more
§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
Compare self to
key and return true if they are equal.