Enum CoopMma 

pub enum CoopMma {
    Fill {
        value: Variable,
    },
    Load {
        value: Variable,
        stride: Variable,
        offset: Variable,
        layout: Option<MatrixLayout>,
    },
    Execute {
        mat_a: Variable,
        mat_b: Variable,
        mat_c: Variable,
    },
    Store {
        mat: Variable,
        stride: Variable,
        offset: Variable,
        layout: MatrixLayout,
    },
    Cast {
        input: Variable,
    },
    RowIndex {
        lane_id: Variable,
        i: Variable,
        matrix: Matrix,
    },
    ColIndex {
        lane_id: Variable,
        i: Variable,
        matrix: Matrix,
    },
    LoadMatrix {
        buffer: Variable,
        offset: Variable,
        vector_size: Option<usize>,
        factor: usize,
        transpose: bool,
    },
    StoreMatrix {
        offset: Variable,
        vector_size: Option<usize>,
        registers: Variable,
        factor: usize,
        transpose: bool,
    },
    ExecuteManual {
        matrix: Matrix,
        registers_a: Variable,
        registers_b: Variable,
        registers_c: Variable,
    },
    ExecuteScaled {
        matrix: Matrix,
        registers_a: Variable,
        registers_b: Variable,
        registers_c: Variable,
        scales_a: Variable,
        scales_b: Variable,
        scales_factor: usize,
    },
}

Cooperative Matrix-Multiply and Accumulate Instruction.

Variants

Fill

Fill the matrix with the value.

Fields

value: Variable

Load

Load the value into the matrix given the stride.

Fields

value: Variable

stride: Variable

offset: Variable

layout: Option<MatrixLayout>

Execute

Executes D=A*B+C;

For implementing a matmul, D=C : C+=A*B

Fields

mat_a: Variable

mat_b: Variable

mat_c: Variable

Store

Store the matrix in an output variable following the stride and the layout.

Fields

mat: Variable

stride: Variable

offset: Variable

layout: MatrixLayout

Cast

Cast a fragment to another type.

Fields

input: Variable

RowIndex

Row index of nth element in the lane

Fields

lane_id: Variable

i: Variable

matrix: Matrix

ColIndex

Column index of nth element in the lane

Fields

lane_id: Variable

i: Variable

matrix: Matrix

LoadMatrix

Execute a CUDA ldmatrix instruction

Fields

buffer: Variable

offset: Variable

vector_size: Option<usize>

factor: usize

transpose: bool

StoreMatrix

Execute a CUDA stmatrix instruction

Fields

offset: Variable

vector_size: Option<usize>

registers: Variable

factor: usize

transpose: bool

ExecuteManual

Manual execute.

Fields

matrix: Matrix

registers_a: Variable

registers_b: Variable

registers_c: Variable

ExecuteScaled

Scaled manual execute.

Fields

matrix: Matrix

registers_a: Variable

registers_b: Variable

registers_c: Variable

scales_a: Variable

scales_b: Variable

scales_factor: usize

Trait Implementations

impl Clone for CoopMma

fn clone(&self) -> CoopMma

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Debug for CoopMma

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

Formats the value using the given formatter.

impl<'de> Deserialize<'de> for CoopMma

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

where __D: Deserializer<'de>,

Deserialize this value from the given Serde deserializer.

impl Display for CoopMma

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

Formats the value using the given formatter.

impl Eq for CoopMma

impl From<<CoopMma as OperationReflect>::OpCode> for OpCode

fn from(value: <CoopMma as OperationReflect>::OpCode) -> OpCode

Converts to this type from the input type.

impl From<CoopMma> for Operation

fn from(value: CoopMma) -> Operation

Converts to this type from the input type.

impl Hash for CoopMma

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 OperationReflect for CoopMma

type OpCode = CmmaOpCode

Type of the op codes for this operation

fn op_code(&self) -> <CoopMma as OperationReflect>::OpCode

Get the opcode for this operation

fn args(&self) -> Option<Vec<Variable>>

Get the list of arguments for this operation. If not all arguments are Variable, returns None instead.

fn from_code_and_args( op_code: <CoopMma as OperationReflect>::OpCode, args: &[Variable], ) -> Option<CoopMma>

Create typed operation from an opcode and a list of arguments. Returns None if not all arguments are Variable.

fn is_commutative(&self) -> bool

Whether this operation is commutative (arguments can be freely reordered). Ignored for single argument operations.

fn is_pure(&self) -> bool

Whether this operation is pure (has no side effects). Things like uniform/plane operations are considered impure, because they affect other units.

impl PartialEq for CoopMma

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

Equality operator ==.

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

Inequality operator !=.

impl Serialize for CoopMma

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 CoopMma

impl TypeHash for CoopMma

fn write_hash(hasher: &mut impl Hasher)

Write the structure of the type to the hasher

fn type_hash() -> u64

Generate a stable hash of the type structure.