Given the `value`, check if itself is an IR value or recursively go through it to check if it contains IR value

Requires an active DSL preprocessor session.

Store a fragment to memory

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

Allocate and initialize barrier memory on the GPU. This function allocates device memory for the barrier. It must be called

Convert a Python primitive value to a Numeric type. :param obj: Python primitive value to convert :type obj: Union[bool, int, float] :ret

Reinterpret the bits of this tensor as a different element type. Total bit width is preserved; the element count adjusts proportionally.

Computes the maximum value from the provided arguments. This function differs from Python's built-in max() in that the return type is determi

Appends an operation and takes ownership

Get the arch enum from the environment variable

< rank of index space

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

Wait for all thread blocks to arrive at the barrier. This method blocks (spins) until all thread blocks have called arrive()

Alias for `div` to match the interface of FastDivmodU64

Unpacks an element from memory

Returns true if element is in a defined state

Returns whether access is valid or not

Obtains a Coord<2> from GemmCoord

Adds a pointer offset in units of Element

Create a new Producer instance with the same state.

Returns whether the block assigned this work should compute the epilogue for the corresponding output tile. For the case of stream-K, this should only

This function is used to check if the expression is a python value. If the expression is a python value, return the boolean value of the expr

Issue Tma Descriptor Prefetch -- ideally from a single thread for best performance

Return the rank of the tensor.

Verify the results

Enhance an IntEnum class with automatic string-to-enum conversion. Builds a reverse mapping from __str__() output to enum members and adds

Accesses the tensor reference pointing to data

//////////////////////////////////////////////////////////////////////////////////////////// Converts a cuDNN status to cutlass::Status

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

Overrides the internal iteration index

Functionally required for serial reduction in the epilogue

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

Constructs a matrix from a uniform element

Arrive at the barrier without waiting. This signals that the calling thread block has reached the barrier point, but does no

Returns the relative dim3 block rank local to the cluster.