Enum LdMatrixForm 

pub enum LdMatrixForm {
    M8N8B16,
    M8N16B4x16P64,
    M8N16B6x16P32,
    M16N16B8,
    M16N16B4x16P64,
    M16N16B6x16P32,
}

The shape of each matrix an ldmatrix loads, and how its elements are laid out in shared memory and in the 32-bit registers they land in. Every row is 16 bytes of shared memory whose address one unit supplies, eight units per 8 rows.

The B4x16P64 and B6x16P32 forms read sixteen packed 4-bit (6-bit) values followed by 64 (32) bits of padding per row, and widen each value into the low bits of its own byte in the registers. An e2m1 operand of the mma kinds that take 8-bit containers sits in bits 2 to 5 of its byte instead, so a 4-bit load must be shifted left by two before it feeds one.

Variants

M8N8B16

8x8 16-bit elements, transposition optional.

M8N16B4x16P64

8x16 4-bit elements widened to bytes, never transposed.

M8N16B6x16P32

8x16 6-bit elements widened to bytes, never transposed.

M16N16B8

16x16 8-bit elements, always transposed.

M16N16B4x16P64

16x16 4-bit elements widened to bytes, always transposed.

M16N16B6x16P32

16x16 6-bit elements widened to bytes, always transposed.

Implementations

impl LdMatrixForm

pub fn registers_per_matrix(self) -> usize

The 32-bit registers each unit receives per matrix.

pub fn transposition(self) -> MatrixMoveTransposition

Whether this form may, must, or cannot transpose.

pub fn admits_matrix_count(self, num_matrices: usize) -> bool

Whether one instruction of this form can load num_matrices matrices.

Trait Implementations

impl Clone for LdMatrixForm

fn clone(&self) -> LdMatrixForm

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Copy for LdMatrixForm

impl Debug for LdMatrixForm

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

Formats the value using the given formatter.

impl<'de> Deserialize<'de> for LdMatrixForm

fn deserialize<__D>( __deserializer: __D, ) -> Result<LdMatrixForm, <__D as Deserializer<'de>>::Error>

where __D: Deserializer<'de>,

Deserialize this value from the given Serde deserializer.

impl Eq for LdMatrixForm

impl Hash for LdMatrixForm

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 Ord for LdMatrixForm

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

fn clamp_to<R>(self, range: R) -> Self

where Self: Sized, R: ClampBounds<Self>,

🔬This is a nightly-only experimental API. (clamp_to)

Restrict a value to a certain range.

impl Parsable for LdMatrixForm

type Arg = ()

Type of the argument that must be passed to the parser.

type Parsed = LdMatrixForm

The type of the parsed entity.

fn parse<'__pliron_parse>( state_stream: &mut Stream<Stream<Stream<Stream<IteratorStream<CharIterator<'__pliron_parse>>, SourcePosition>>>, State<'__pliron_parse>>, arg: <LdMatrixForm as Parsable>::Arg, ) -> Result<(<LdMatrixForm as Parsable>::Parsed, Commit<()>), Commit<Tracked<<Stream<Stream<Stream<Stream<IteratorStream<CharIterator<'__pliron_parse>>, SourcePosition>>>, State<'__pliron_parse>> as StreamOnce>::Error>>>

Define a parser using existing combinators and call into on Parser::parse_stream to get the final ParseResult. Use state_stream.state as necessary.

fn parser<'a>( arg: Self::Arg, ) -> Box<dyn Parser<Stream<Stream<Stream<Stream<IteratorStream<CharIterator<'a>>, SourcePosition>>>, State<'a>>, PartialState = (), Output = Self::Parsed> + 'a>

where Self::Parsed: 'a,

Get a parser combinator that can work on StateStream as its input.

impl PartialEq for LdMatrixForm

fn eq(&self, other: &LdMatrixForm) -> bool

Equality operator ==.

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

Inequality operator !=.

impl PartialOrd for LdMatrixForm

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

fn fmt( &self, ctx: &Context, state: &State, fmt: &mut Formatter<'_>, ) -> Result<(), Error>

fn disp<'t, 'c>(&'t self, ctx: &'c Context) -> Box<dyn Display + 'c>

where 't: 'c,

Get a Display’able object from the given Context and default State.

fn print<'t, 'c>( &'t self, ctx: &'c Context, state: &State, ) -> Box<dyn Display + 'c>

where 't: 'c,

Get a Display’able object from the given Context and State.

impl Serialize for LdMatrixForm

fn serialize<__S>( &self, __serializer: __S, ) -> Result<<__S as Serializer>::Ok, <__S as Serializer>::Error>

where __S: Serializer,

Serialize this value into the given Serde serializer.

impl StructuralPartialEq for LdMatrixForm