Argument wrapper for GEMM in CUTLASS 2 or 3. It returns either 2x arguments or 3x arguments depending on the `arch` field specified in `opera

Constructs the function object, possibly loading from pointers in host memory

< Layout function

< Layout function

Extract MLIR values from this object. :return: List of MLIR values representing this object's data :rtype: List[ir.Value]

Post-initialization validation for TMEM load operations. Performs comprehensive validation of operation parameters and architecture

Helper method to convert an arguments list to its corresponding MLIR types. This method converts numeric meta types and types convertible vi

Convert a vector or tensor from column-major order to row-major order. :param vec: The input vector in column-major order :type vec: ir.Value

Create FP4 tensor. Args: logical_shape: shape in FP4 elements (packed_dim size must be even). packed_dim: dimension t

LUT x4 _e2m1_to_bf16_x4: Direct E2M1->BF16 (converts 4 E2M1 to 4 BF16 in one call)

Flatten a nested dict/list structure of MLIR values into a flat list. Local copy to avoid circular imports with dsl.py.

Returns the batch stride of ``tensor``. If ``tensor`` is only rank-2, batch stride is 0. :param tensor: tensor object to process

make acc tensor as mn layout view :param acc: input tensor :type acc: cute.Tensor :return: acc tensor mn layout view

Convert region_result to a list if it is not already a list If region_result is a list, return it as is. If region_result is

Prepare a ref tensor: make ``k_dim`` stride-1 and optionally pad. Args: tensor: input data tensor (A or B). k_dim: wh

Set the epilogue functor based on the provided activation function

Compile-only helper: JIT-compiles and runs the kernel once without validation. Uses iterations=1 (not 0) because testing.benchmark divides by ite

Verifies that ``tensor`` has rank greater than 1 :param tensor: object representing a tensor passed in to verify, or ``None`` if no

Provide the accumulator value to the user.

Acquire pipeline stage and synchronize all warps. TMA warp waits for previous TMA operation to the same stage to complete (allowing

Adds a pointer offset in units of element

Adds a pointer offset in units of Element

Adds a pointer offset in units of Element

Accesses a bit within the predicate vector.

Called at the start of one step before starting accumulator exchange

Cudnn compute type seems to be hardcoded to float (To handle a possible cudnn issue)

Clears the predicates

Clears the predicate set efficiently

Fence SMEM writes and synchronize all warps. All warps fence their SMEM writes to make them visible to consumer All warps th

Performs an expected transaction bytes decrement without doing an arrive operation

Computes predicates based on internally tracked per-thread offset.

Compute the assembled (swizzled 32_4_4) scale tensor shape. Swizzle 32_4_4 pads each group's scale to rows=round_up(non_K, 128), cols=ro

Tries to wait for a consumer stage, non-blocking.

The basic operation kind is prefixed with a letter indicating the accumulation type.

Return size of codomain of layout or tensor. Return static value if type is static. For a layout ``L = S:D`` where ``S`` is the shape and ``D`` i

Create one raw, non-swizzled scale tensor with exact values in {1, 2}.

Dequantize a single 2D tensor using raw block scales into fp32.

Computes integer division and modulus using precomputed values. This is computationally inexpensive. Simply returns the quotient

Sets all predicates to true

Clears the predicate set efficiently

Partitions source and destination tensors for a TMA store or SIMT store.

Make tiledCopy for shared memory store, then use it to partition register array (source) and shared memory (destination). :p

Make tiledCopy for tensor memory load, then use it to partition tensor memory (source) and register array (destination). :pa

Make tiledCopy for tensor memory load, then use it to partition tensor memory (source) and register array (destination). :param gemm_kernel:

Given the `obj`, recursively go through it to extract all contained IR attributes as list of MLIR attributes. This is used for generating ker

Find the runtime libraries that needs to be available for loading modules. :param enable_tvm_ffi: Whether to enable TVM-FFI. :type enabl

Execute the given function for each supplemental written tensor.

Generate all TensorOp math instructions for FP8 MMA that are supported by SM100 at or above the given level. Args: level: The global