Adds a pointer offset in units of Element

Load addend source fragment from global memory

Returns whether access is valid or not

Returns a reference to the element at a given Coord

Issuing the griddepcontrol.wait instruction enforces no global memory access prior to this istruction. This ensures the correctness of global memory a

Notify the iterator which k-group it is currently pointing to. This does not advance the iterator. Rather, it overrides its internal tracking with co

Replaces the managed object, deleting the old object.

Establishes an ordering w.r.t previously issued cp.async instructions. Does not block.

Converts to float

Returns stride as MatrixCoord

Replace tensor pointer by bringing the tensormap from GMEM into the shared memory

Returns a reference to the element at a given Coord

Stores a fragment to memory

Static method that returns the packed logical size of any TensorExtent

Marks the commit point for one or more sized batch of warpgroup MMAs.

Advances an iterator along logical dimensions of matrix in units of whole threadblock tiles

Returns a TensorRef pointing to the first element of the tensor.

Accesses the x element of the imaginary part of the quaternion

Accesses the y element of the imaginary part of the quaternion

Accesses the z element of the imaginary part of the quaternion

Returns a pointer to the vector starting at the current coordinate

Returns the GEMM volume in thread block tiles

Perform a Producer Epilogue to prevent early exit of ctas in a Cluster

Bitcasts to CUDA's half type

Exit of the subtile load loop.

Loads a fragment from memory

Loads a fragment from memory

Returns the smallest value in the half-open range [a, a+b) that is a multiple of b

Adds a pointer offset in units of Element

Store a fragment to memory

Replace tensor dims and strides for GEMMs by bringing the tensormap from GMEM into the shared memory

Get 1D ctaid in a cluster.

Issue a shared memory fence for async operations

Maximum finite value

Stores a fragment to memory at the location pointed to by the iterator

Alias for `div` to match the interface of FastDivmodU64

Unpacks an element from memory