Compute the permuted higher-rank tensor shape from the original shape.

Prevents early exit of producer blocks in Cluster. This should be called once before kernel exits.

< Layout function

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

Create an i32 constant with the given value.

Given the `obj`, recursively go through it to extract all contained IR values. Args: obj: The object to extract MLIR values from

Implemented by subclasses.

Returns the extent of the tensor

Returns the layout object

< Efficiently disables all accesses guarded by mask

Returns the ceiling of (a / b)

Returns the layout object

Copies data from device to host

Returns the logical capacity in terms of number of elements. May be larger than the size().

Gets pointer to host data

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

Converts a ExecutionMode enumerant to a string

Returns a const iterator

Returns a pointer to the managed object or nullptr if no object is owned.

This is the method we expect to use going forward

Returns the rank of the tensor

Gets runtime fields of the Atom. Some Atoms have runtime state, for example a tcgen05 MMA Atom .. code-block:: python

Create a new python object by populating containing MLIR values with new values. Args: obj: The original object to use as a template

Writes the entire listing to a file

Arrive at the barrier AND wait for all other thread blocks. This is the most common barrier operation - it combines arrive()

Overrides the internal iteration index

Helper function to log messages. Change logger.info to another level here if needed.

Prevents early exit of producer blocks in Cluster. This should be called once before kernel exits.

Allocate a tensor in shared memory. Note: Currently only supports static layouts. Dynamic layouts are not supported. :param element_

Appends an entry

Barrier arrive on local smem

The pass that is run through the self.dag_ir

Atomically wait for data and advance to next pipeline stage. This is a convenience method that combines wait() and advance() into a single

< Efficiently disables all accesses guarded by mask

Computes the offset of an index from the origin of the tensor

Acquire the current buffer and advance to the next pipeline stage. This method combines the acquire() and advance() operations into a single

Overrides the internal iteration index

Returns a warp-uniform value indicating the canonical warp index of the calling threads. Threads within the warp must be converged.