A range-like object for dynamic loop iteration in the DSL. This class provides a range interface similar to Python's built-in range, but

Raised when an error occurs during JIT-time code generation in the DSL.

MatrixCoord wraps Coord<2, int> to provide a helper for accessing named dimensions. Classes expecting a coordinate in the rank=2 index space of a matr

Coordinate in pitch-linear space

An exception class for Op construction errors.

Utility class to map transpose and non-transpose terminology to row- and column-major terminology

Boolean type representation in the DSL. This class represents boolean values in the DSL, with a width of 1 bit. It supports conversion from v

IEEE half-precision floating-point type

Floating-point type with 8 bits of exponent and 7 bits of mantissa.

GemmCoord is a structure derived from Coord<3> that specifies a location within the coordinate space of a GEMM problem.

Structure representing an element in source memory

Tensor Float 32 data type

Object to encapsulate the fast division+modulus operation. This object precomputes two values used to accelerate the computation and is best used whe

Mapping function for pitch-linear memory

Epilogue Visitor Testbed

Structure describing the properties of a tensor

Utility class to assert that an exception was raised when expected Example: .. highlight:: python .. code-block:: python w

A class representing thread local data from CuTe Tensor in value semantic and immutable. :param value: Flatten vector as ir.Value holding logic d

Object to encapsulate the fast division+modulus operation for 64b integer division. This object precomputes two values used to accelerate the computa

////////////////////////////////////////////////////////// floating-point 8 type : E4M3 //////////////////////////////////////////////////////////

Forward declaration of float_e5m2_t to define float_e4m3_t <=> float_e5m2_t conversions in class float_e4m3_t

Mask object

Container for TMA Copy Atom and related data. This class uses software composition to bundle a CopyAtom with the SMEM layout and TMA ten

The maximum swizzle size to use This class, like Splits above makes it harder to confuse the order of arguments of the various run(...) functions in

Defines a canonical 4D coordinate used by tensor operations.

Defines a canonical 5D coordinate used by tensor operations.

Wrap an MLIR ``vector<NxTy>`` register value with DSL type information. Provides element extraction (``vec[i]`` / ``vec[a:b]``), element-wise

Pipeline for tile scheduling (using CLC) to all warps.

Raised when an error occurs during AST preprocessing or visiting in the DSL.

Internal representation of constrained integer types with divisibility information. IntValue serves as a proxy for constrained integer types in t

! Unsigned 128b integer type

RAII CUDA buffer container

Raised when a feature of the DSL is not implemented yet.

Object to encapsulate the fast division+modulus operation for 64b integer division in which the divisor is a power of two.

A deterministic set implementation for ordered operations.

A class to represent information about a work tile. :ivar tile_idx: The index of the tile. :type tile_idx: cute.Coord :ivar is_valid_tile

BatchedGemmCoord is a structure derived from Coord<4> that specifies a location within the coordinate space of a batched GEMM problem.

Context manager for handling regions during AST transformations. This class is used to manage region-scoped state during DSL preprocessing.

The tensor abstracts the data type

variables: pointer, integer, floating-point, boolean, etc. Parameters ---------- name : str The parameter name. dtype : str |

Bulk tensor asynchronous multicast GMEM to SMEM Copy Operation using the TMA unit. TMA multicast operations are issued by a single thread wi

Work tile information for MoE scheduler. Contains CTA-level tile information for executor warps: - expert_idx: Which expert (-1 mean

r"""Builtin Tensor Type as an IR value supporting standard iterator and layout.. :param value: The MLIR operation result value to initialize the

Bulk tensor asynchronous GMEM to SMEM Copy Operation using the TMA unit. TMA copy operations are issued by a single thread within a warp, bu

The universal Copy Operation. This operation is equivalent to the ``a = b`` assignment without any extra memory attributes. For advance

Custom exception for setup errors.

device_id: CUDA device ID to get the hardware info.

Pipeline state contains an index and phase bit corresponding to the current position in the circular buffer.

An enumeration for the number of repetitions of a given TMEM copy within the instruction.

Mapping function for row-major matrices.

8x8 ``stmatrix`` Operation. See the `PTX documentation <https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructio