This is a wrapper around the _cute_ir.make_atom operation, providing default value for the values argument.

Replaces the managed object, deleting the old object.

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

Returns the GEMM volume in thread block tiles

Return size of domain of layout or tensor. Computes the size (number of elements) in the domain of a layout or tensor. For layouts, this corr

Maximum finite value

Signal that data production is complete for the current stage. This allows consumers to start processing the data. :param handle: Op

Returns stride as MatrixCoord

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

Check CUDA errors and provide detailed error messages. :param result: The result of the CUDA operation. :type result: tuple(CUresult, ...)

Create a pointer from a memory address :param dtype: Data type of the pointer elements :type dtype: Type[Numeric] :param value: Memory ad

Prefetch a TMA descriptor from global memory. This function prefetches the TMA descriptor pointed to by tma_desc_ptr into the TMA descri

Final processing step for FMHA that computes log-sum-exp (LSE) and scales the output. This function performs the following operation

Get the maximum number of active clusters for a given cluster size. When a stream from a green context is provided, the occupancy ca

Stores a fragment to memory

Returns a reference to the element at a given Coord

Get 1D ctaid in a cluster.

Create an i64 constant with the given value.

Static method that returns the packed logical size of any TensorExtent

Check if a value is statically known at compile time. In CuTe, static values are those whose values are known at compile time, as opposed to

Create a CuTe Layout object from shape and optional stride information. A Layout in CuTe represents the mapping between logical and physical coor

Bitcasts to CUDA's half type

Returns a pointer

Helper to add an offset in number of elements, assuming this offset is divisible by the vector size.

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

Try to get the JitCompiledFunction object for the given key. If the key is not in the cache, None will be returned. This has the same

Get the meta data of the node

Syncs the threadblock or cluster

Compute element-wise square root of the input tensor. :param a: Input tensor :type a: Union[TensorSSA, Numeric] :param fastmath: Enable f

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

Allocate a block of tensor memory. This method allocates a block of tensor memory from allocator warp and returns a handle to retrieve

Returns a pointer to the vector starting at the current coordinate

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

Exit of the subtile load loop.

Loads a fragment from memory

< Layout function

Get the barrier pointer for the current pipeline stage. :return: Pointer to the barrier for the current stage :rtype: cute.Po

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

Filtering operations based on various criteria

Deallocate the tensor memory. This method sync on mbarrier (for two cta use case) and deallocates the tensor memory from the allocator warp.

Fences the mbarrier_init and sends an arrive if using clusters.

Helper function to check CUDA API errors.

Copy of :meth:`ast.NodeTransformer.generic_visit` with support for inserting statements during expression visits. This version provi

Loads a fragment from memory

Relinquish the tensor memory allocation permit. This method relinquishes the tensor memory allocation permit for the allocator warp, promisin

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

Returns the stride of the layout

Adds a pointer offset in units of Element