### Static Persistent Kernel Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/blackwell_cluster_launch_control.html

Shows the implementation of a static persistent kernel. It involves setup, fetching work iteratively using a static tile scheduler, and performing computations until no valid work remains.

```cpp
// Static Persistent Kernel
__device__ static_persistent_kernel(...) {
  setup_common_data_structures(...);
  dim3 workCoordinates = blockIdx;
  bool isValidId;
  do {
    coordinate_specific_compute(workCoordinates);
    std::tie(isValidId, workCoordinates) = staticTileScheduler.fetch_next_work();
  } while (isValidId);
}


```

--------------------------------

### Asynchronous Pipeline Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/pipeline.html

Demonstrates the setup and usage of a 4-stage asynchronous pipeline with 2 producer threads and 1 consumer thread. Use this to manage data dependencies and synchronization between threads performing different roles in a computation pipeline.

```cpp
// 4-stage Pipeline
static constexpr int NumStages = 4;
using MainloopPipeline = typename cutlass::PipelineAsync<NumStages>;
using PipelineState = typename cutlass::PipelineState<NumStages>;

// 2 producer threads and 1 consumer thread 
typename MainloopPipeline::Params params;
params.producer_arv_count = 2;
params.consumer_arv_count = 1;
MainloopPipeline pipeline(shared_storage.storage, params);
  
// Producer threads
if (thread_idx == 0 or thread_idx == 1) {
  PipelineState smem_pipe_write = cutlass::make_producer_start_state<MainloopPipeline>();
  for ( ; iter > 0; --iter) {
    pipeline.producer_acquire(smem_pipe_write);

    // Producer ops
    // If any memory operations are involved, then we also need
    // to guarantee that writes are completed and visible to consumer(s).

    pipeline.producer_commit(smem_pipe_write);
    ++smem_pipe_write;
  }
}
else if (thread_idx == 2) {
  PipelineState smem_pipe_read;
  for (; iter > 0; --iter) {
    pipeline.consumer_wait(smem_pipe_read);

    // Consumer ops

    pipeline.consumer_release(smem_pipe_read);
    ++smem_pipe_read;
  }
}

```

--------------------------------

### Blackwell Dynamic Persistent Kernel Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/blackwell_cluster_launch_control.html

Demonstrates a dynamic persistent kernel for Blackwell, utilizing a dynamic tile scheduler. It includes setup, fetching the next work item, and continuing computation as long as valid work IDs are available.

```cpp
// Dynamic Persistent Kernel
__device__ clc_dynamic_persistent_kernel(...) {
  setup_common_data_structures(...);
  dim3 workCoordinates = blockIdx;
  dim3 newClcID;
  bool isValidId;
  do {
    coordinate_specific_compute(workCoordinates);
    std::tie(isValidId, newClcID) = clcTileScheduler.fetch_next_work();
    workCoordinates = newClcID;
  } while (isValidId);
}


```

--------------------------------

### Build CUTLASS Example with synclog Enabled

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/utilities.html

After enabling `synclog` via CMake, navigate to your desired CUTLASS example directory and use `make` to build it. This ensures the example is compiled with `synclog` support.

```bash
$ cd examples/54_hopper_fp8_warp_specialized_gemm
$ make

```

--------------------------------

### Run Example and Capture synclog Output

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/utilities.html

Execute the compiled CUTLASS example and redirect its output to a file (e.g., `synclog.txt`). Set the `--iterations=0` flag to ensure `synclog` information is captured only for the reference run, simplifying analysis.

```bash
$ ./54_hopper_fp8_warp_specialized_gemm --iterations=0 &> synclog.txt

```

--------------------------------

### CUTLASS SDK Examples Directory Structure

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/code_organization.html

This snippet outlines the directory structure for the CUTLASS SDK examples. Each subdirectory demonstrates a specific CUTLASS component or computation, such as basic GEMM, utilities, iterators, and fused operations.

```text
examples/
  00_basic_gemm/             # launches a basic GEMM with single precision inputs and outputs

  01_cutlass_utilities/      # demonstrates CUTLASS Utilities for allocating and initializing tensors
  
  02_dump_reg_smem/          # debugging utilities for printing register and shared memory contents
  
  03_visualize_layout/       # utility for visualizing all layout functions in CUTLASS

  04_tile_iterator/          # example demonstrating an iterator over tiles in memory

  05_batched_gemm/           # example demonstrating CUTLASS's batched strided GEMM operation

  06_splitK_gemm/            # exmaple demonstrating CUTLASS's Split-K parallel reduction kernel

  07_volta_tensorop_gemm/    # example demonstrating mixed precision GEMM using Volta Tensor Cores

  08_turing_tensorop_gemm/   # example demonstrating integer GEMM using Turing Tensor Cores

  10_planar_complex/         # example demonstrating planar complex GEMM kernels

  11_planar_complex_array/   # example demonstrating planar complex kernels with batch-specific problem sizes

  12_gemm_bias_relu/         # example demonstrating GEMM fused with bias and relu activation function

  13_fused_two_gemms/        # example demonstrating two GEMMs fused into one kernel

```

--------------------------------

### Non-persistent Kernel Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/blackwell_cluster_launch_control.html

Illustrates the structure of a non-persistent kernel in the CLC programming model. This kernel performs setup and then coordinate-specific computation.

```cpp
// Non-persistent kernel
__device__ non_persistent_kernel(...) {
  setup_common_data_structures();
  dim3 workCoordinates = blockIdx;
  coordinate_specific_compute(workCoordinates);
}


```

--------------------------------

### Pipeline Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/pipeline.html

An example demonstrating the usage of asynchronous pipeline classes with multiple producer and consumer threads for a 4-stage pipeline.

```APIDOC
## Pipeline Example

### Description
This example illustrates the setup and usage of an asynchronous pipeline with 2 producer threads and 1 consumer thread across 4 stages. It shows how `producer_acquire`, `producer_commit`, `consumer_wait`, and `consumer_release` are used in a loop.

### Code
```cpp
// 4-stage Pipeline
static constexpr int NumStages = 4;
using MainloopPipeline = typename cutlass::PipelineAsync<NumStages>;
using PipelineState = typename cutlass::PipelineState<NumStages>;

// 2 producer threads and 1 consumer thread 
typename MainloopPipeline::Params params;
params.producer_arv_count = 2;
params.consumer_arv_count = 1;
MainloopPipeline pipeline(shared_storage.storage, params);
  
// Producer threads
if (thread_idx == 0 or thread_idx == 1) {
  PipelineState smem_pipe_write = cutlass::make_producer_start_state<MainloopPipeline>();
  for ( ; iter > 0; --iter) {
    pipeline.producer_acquire(smem_pipe_write);

    // Producer ops
    // If any memory operations are involved, then we also need
    // to guarantee that writes are completed and visible to consumer(s).

    pipeline.producer_commit(smem_pipe_write);
    ++smem_pipe_write;
  }
}
else if (thread_idx == 2) {
  PipelineState smem_pipe_read;
  for (; iter > 0; --iter) {
    pipeline.consumer_wait(smem_pipe_read);

    // Consumer ops

    pipeline.consumer_release(smem_pipe_read);
    ++smem_pipe_read;
  }
}
```

### Notes
This is a basic example. Different pipeline classes and configurations are available. Refer to CUTLASS unit tests for more advanced usage.
```

--------------------------------

### Install nvidia-matmul-heuristics

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/heuristics.html

Install the necessary Python package for GEMM heuristics using pip. This is the recommended installation method.

```bash
pip install nvidia-matmul-heuristics

```

--------------------------------

### Install TVM FFI

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/compile_with_tvm_ffi.html

Install the apache-tvm-ffi package and optionally the torch-c-dlpack-ext for better PyTorch tensor performance.

```bash
pip install apache-tvm-ffi
# optional package for improved torch tensor calling performance
pip install torch-c-dlpack-ext

```

--------------------------------

### Tracing Mode Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/dsl_code_generation.html

Illustrates the use of tracing mode with `preprocess=False` for kernels guaranteed to be straight-line arithmetic.

```python
from cutlass_python import jit

@jit(preprocess=False)
def my_kernel(a, b, c):
    # Straight-line arithmetic operations
    c[:] = a[:] + b[:]

```

--------------------------------

### Producer-Consumer Pipeline Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/pipeline.html

Demonstrates the basic usage of PipelineProducer and PipelineConsumer for asynchronous data production and consumption. Shows how to acquire, produce, and commit data within a pipeline loop.

```python
pipeline = PipelineAsync.create(...)
producer, consumer = pipeline.make_participants()
for i in range(iterations):
    # Try to acquire the current buffer without blocking
    try_acquire_token = producer.try_acquire()

    # Do something else independently
    ...

    # Wait for current buffer to be empty & Move index to next stage
    # If try_acquire_token is True, return immediately
    # If try_acquire_token is False, block until buffer is empty
    handle = producer.acquire_and_advance(try_acquire_token)

    # Produce data
    handle.commit()

```

--------------------------------

### Example of zipped_divide

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Demonstrates how to use zipped_divide to partition a layout with a given tiler. The result shape is shown.

```python
layout = cute.make_layout((128, 64), stride=(64, 1))
tiler = (8, 8)
result = cute.zipped_divide(layout, tiler)  # result shape: ((8, 8), (16, 8))
```

--------------------------------

### Device-wide GEMM API Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/gemm_api.html

Launches a mixed-precision GEMM operation targeting Volta Tensor Cores. Ensure necessary types and architectures are included in your project.

```cpp
  using Gemm = cutlass::gemm::device::Gemm<
    cutlass::half_t,                           // ElementA
    cutlass::layout::ColumnMajor,              // LayoutA
    cutlass::half_t,                           // ElementB
    cutlass::layout::ColumnMajor,              // LayoutB
    cutlass::half_t,                           // ElementOutput
    cutlass::layout::ColumnMajor,              // LayoutOutput
    float,                                     // ElementAccumulator
    cutlass::arch::OpClassTensorOp,            // tag indicating Tensor Cores
    cutlass::arch::Sm70                        // tag indicating target GPU compute architecture
  >;

  Gemm gemm_op;
  cutlass::Status status;
 
  //
  // Launch GEMM on the device
  //
 
  status = gemm_op({
    {m, n, k},
    {ptrA, lda},
    {ptrB, ldb},
    {ptrC, ldc},
    {ptrD, ldd},
    {alpha, beta}
  });

  if (status != cutlass::Status::kSuccess) {
    return -1;
  }

```

--------------------------------

### Preprocessor Mode Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/dsl_code_generation.html

Demonstrates the default preprocessor mode (`preprocess=True`) which combines AST rewrite and tracing for handling loops and branches.

```python
from cutlass_python import jit

@jit(preprocess=True)
def my_kernel_with_control_flow(a, b, c):
    # Kernel with potential loops and branches
    for i in range(a.shape[0]):
        if a[i] > 0:
            c[i] = a[i] * b[i]
        else:
            c[i] = a[i] + b[i]

```

--------------------------------

### Macro for Host and Device Functions

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/programming_guidelines.html

Macros should only be used when the preprocessor is essential. This example shows a macro for functions that can run on both the host and the device.

```cpp
#define CUTLASS_MACROS_USE_ALL_CAPS inline __host__ __device__
```

--------------------------------

### Tensor Reduction Examples

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Demonstrates tensor reduction operations with different reduction profiles. Specifies which dimensions to reduce, keeping others.

```python
reduce(f32 o (4,)) 
  => f32
```

```python
reduce(f32 o (4, 5)) 
  => f32
```

```python
reduce(f32 o (4, (5, 4)), reduction_profile=(None, 1)) 
  => f32 o (4,)
```

```python
reduce(f32 o (4, (5, 4)), reduction_profile=(None, (None, 1))) 
  => f32 o (4, (5,))
```

--------------------------------

### Column-Major Layout Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/layout.html

Demonstrates how CUTLASS's `layout::ColumnMajor` represents a Fortran-style column-major matrix layout, commonly used in BLAS.

```cpp
layout::ColumnMajor layout(lda); 

int offset = layout({row, column});     // returns row  + lda * column

```

--------------------------------

### GEMM with Sequential Input Distribution

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/profiler.html

Use a sequential distribution for input values in GEMM operations, specifying the start value and the increment.

```bash
cutlass_profiler --operation=Gemm --dist=sequential,start:0,delta:1
```

--------------------------------

### Specify GEMM K dimension range

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/profiler.html

Use ranges like `start:end:increment` or `start:end` for integer arguments to sweep across values. This example sweeps the K dimension from 8 to 4096 in increments of 8.

```bash
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sgemm_128x128_nn --m=4352 --n=4096 --k=8:4096:8

```

--------------------------------

### Getting the Data Pointer of a Memory Range

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Illustrates how to obtain the starting memory address (pointer) for a given memory range.

```python
data_ptr()
```

--------------------------------

### Create and Use PipelineAsync

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/pipeline.html

Demonstrates how to create a PipelineAsync with a specified number of stages and producer/consumer groups, and then use the producer and consumer participants to write and read data.

```python
# Create pipeline with 5 stages
pipeline = PipelineAsync.create(
    num_stages=5,                   # number of pipeline stages
    producer_group=producer_warp,
    consumer_group=consumer_warp
    barrier_storage=smem_ptr,       # smem pointer for array of mbarriers in shared memory
)

producer, consumer = pipeline.make_participants()
# Producer side
for i in range(num_iterations):
    handle = producer.acquire_and_advance()  # Wait for buffer to be empty & Move index to next stage
    # Write data to pipeline buffer
    handle.commit()   # Signal buffer is full

# Consumer side
for i in range(num_iterations):
    handle = consumer.wait_and_advance()     # Wait for buffer to be full & Move index to next stage
    # Read data from pipeline buffer
    handle.release()  # Signal buffer is empty

```

--------------------------------

### Compile and Run Kernel with Variadic Tuple

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/compile_with_tvm_ffi.html

Demonstrates compiling a kernel that accepts a variadic tuple for flexible arguments. Shows how to compile with and without extra values and execute the compiled functions.

```python
import cutlass
import torch
from cutlass import cute

@cute.kernel
def device_add_one(a: cute.Tensor, b: cute.Tensor, extra_value: tuple):
   threads_per_block = 128
   cta_x_, _, _ = cute.arch.block_idx()
   tid_x, _, _ = cute.arch.thread_idx()
   tid = cta_x_ * threads_per_block + tid_x
   if tid < a.shape[0]:
      if cutlass.const_expr(len(extra_value) != 0):
            b[tid] = a[tid] + 1 + extra_value[0]
      else:
            b[tid] = a[tid] + 1

@cute.jit
def add_one_with_extra_value(a: cute.Tensor, b: cute.Tensor, extra_value: tuple):
   n = a.shape[0]
   threads_per_block = 128
   blocks = (n + threads_per_block - 1) // threads_per_block
   device_add_one(a, b, extra_value).launch(grid=(blocks, 1, 1), block=(threads_per_block, 1, 1))

def example_add_one_with_variadic_tuple():
   n = cute.sym_int()
   a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   compiled_add_one_no_extra = cute.compile(
      add_one_with_extra_value, a_cute, b_cute, (),
      options="--enable-tvm-ffi"
   )
   compiled_add_one_with_extra = cute.compile(
      add_one_with_extra_value, a_cute, b_cute, (cute.Float32(4),),
      options="--enable-tvm-ffi"
   )
   a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
   b_torch = torch.empty(10, dtype=torch.float32, device="cuda")
   compiled_add_one_no_extra(a_torch, b_torch, ())
   print("result of b_torch after compiled_add_one_no_extra(a_torch, b_torch, ())")
   print(b_torch)
   compiled_add_one_with_extra(a_torch, b_torch, (4,))
   print("result of b_torch after compiled_add_one_with_extra(a_torch, b_torch, (4,))")
   print(b_torch)

example_add_one_with_variadic_tuple()

```

--------------------------------

### Get CUDA Error Enum

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils.html

Internal utility to get the CUDA error enumeration string.

```python
def _cudaGetErrorEnum(_error : Any_) -> str:
    pass
```

--------------------------------

### Registering JIT Argument Adapter for 3rd Party Types

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/dsl_jit_arg_generation.html

Demonstrates how to use an adaptor class to bridge CuTe DSL with a 3rd party framework type, enabling JIT argument generation without modifying the original type.

```python
@cutlass.register_jit_arg_adapter(MyFrameworkObject)
class MyFrameworkObjectAdapter:
    """
    Convert a 3rd party framework object to a JIT function argument with JitArgument protocol
    """

    def __init__(self, arg):
        self._arg = arg

    def __c_pointers__(self):
        # Convert the framework object to a C-ABI compatible object
        # thru its C-ABI interface
        return [self._arg.get_cabi_pointer()]

    def __get_mlir_types__(self):
        # Return the list of MLIR types the framework object represents
        return [self._arg.get_data().mlir_type]

    def __new_from_mlir_values__(self, values):
        # Convert the MLIR values back to the framework object
        return MyFrameworkObject(values[0])

```

--------------------------------

### Barrier Setup with elect_one

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute_nvgpu_cpasync.html

Barrier initialization and transaction byte setup for TMA operations require `elect_one()`. The actual TMA copy does not.

```python
# Barrier setup requires elect_one
with cute.arch.elect_one():
    cute.arch.mbarrier_init(barrier_ptr, arrival_count)
    cute.arch.mbarrier_expect_tx(barrier_ptr, num_tma_bytes)

# TMA copy does NOT need elect_one
cute.copy(tma_atom, gmem_tensor, smem_tensor, tma_bar_ptr=barrier_ptr)
```

--------------------------------

### Get Runtime Fields of an Atom

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Use the `get` method to access runtime state of an Atom, such as the ACCUMULATE field for a tcgen05 MMA Atom.

```python
tiled_mma = cute.make_tiled_mma(some_tcgen05_mma_op)
accum = tiled_mma.get(cute.nvgpu.tcgen05.Field.ACCUMULATE)
```

--------------------------------

### make_tiled_copy_tv

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Create a tiled copy given separate thread and value layouts. A TV partitioner is inferred based on the input layouts. The input thread layout must be compact.

```APIDOC
## make_tiled_copy_tv(_atom: CopyAtom, _thr_layout: Layout, _val_layout: Layout, *_, _loc: Location | None = None, _ip: InsertionPoint | None = None) -> TiledCopy

### Description
Create a tiled copy given separate thread and value layouts. A TV partitioner is inferred based on the input layouts. The input thread layout must be compact.

### Parameters

#### Path Parameters
* **atom** (CopyAtom) - Copy atom.
* **thr_layout** (Layout) - Layout mapping from (TileM,TileN) coordinates to thread IDs (must be compact).
* **val_layout** (Layout) - Layout mapping from (ValueM,ValueN) coordinates to value IDs.

#### Optional Parameters
* **loc** (Location | None) - Source location for MLIR, defaults to None.
* **ip** (InsertionPoint | None) - Insertion point, defaults to None.

### Returns
A tiled copy for the partitioner.

### Return type
TiledCopy
```

--------------------------------

### MMA_Traits Specialization Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/0t_mma_atom.html

An example specialization of MMA_Traits for the SM70_8x8x4_F32F16F16F32_NT Operation. It defines logical compute types, shape, thread ID mapping, and layout for matrices A, B, and C.

```cpp
template <>
struct MMA_Traits<SM70_8x8x4_F32F16F16F32_NT>
{
  using ValTypeD = float;
  using ValTypeA = half_t;
  using ValTypeB = half_t;
  using ValTypeC = float;

  using Shape_MNK = Shape<_8,_8,_4>;
  using ThrID   = SM70_QuadPair;
  using ALayout = SM70_8x4_Col;
  using BLayout = SM70_8x4_Col;
  using CLayout = SM70_8x8_32b;
};


```

--------------------------------

### CUTLASS 3.0 GEMM procedural name example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/profiler.html

This example shows a CUTLASS 3.0 GEMM procedural name, illustrating the new naming convention for kernels targeting NVIDIA Hopper architecture and beyond.

```plaintext
cutlass3x_sm90_tensorop_gemm_f16_f16_f32_f16_f32_{optional-mixed-dtype-config}_128x128x64_2x1x1_0_ntn_align8

```

--------------------------------

### Manually Create SMEM Layouts

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/mma_docs/tcgen05_programming.html

Demonstrates manual creation of SMEM layouts by defining a swizzle, tiling to the MMA shape, and specifying the order for contiguous dimensions, without using utility functions.

```python
swizzle = cute.Swizzle(3, 4, 3)
mma_tile = tiled_mma.partition_shape_A((mma_tiler_mnk[0], mma_tiler_mnk[2]))
smem_tile = tcgen05.tile_to_mma_shape(swizzle, mma_tile, order=(1, 2, 3))

```

--------------------------------

### Tensor Initialization and Access Examples

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/03_tensor.html

Illustrates creating Tensors with specified shapes and strides, and accessing/modifying elements using operator[] with natural coordinates, variadic operator(), and array-like access.

```cpp
Tensor A = make_tensor<float>(Shape <Shape < _4,_5>,Int<13>>{}, 
                              Stride<Stride<_12,_1>,    _64>{});
float* b_ptr = ...;
Tensor B = make_tensor(b_ptr, make_shape(13, 20));

// Fill A via natural coordinates op[]
for (int m0 = 0; m0 < size<0,0>(A); ++m0)
  for (int m1 = 0; m1 < size<0,1>(A); ++m1)
    for (int n = 0; n < size<1>(A); ++n)
      A[make_coord(make_coord(m0,m1),n)] = n + 2 * m0;

// Transpose A into B using variadic op()
for (int m = 0; m < size<0>(A); ++m)
  for (int n = 0; n < size<1>(A); ++n)
    B(n,m) = A(m,n);

// Copy B to A as if they are arrays
for (int i = 0; i < A.size(); ++i)
  A[i] = B[i];

```

--------------------------------

### Blackwell GEMM Kernel Builder Setup

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/blackwell_functionality.html

Skeleton C++ code for setting up a block-scaled GEMM kernel using the CUTLASS collective builder. This includes defining tensor element types, layouts, alignment, and performance parameters like tile and cluster shapes.

```cpp
  ///////////////////////////////////////////////////////////
  //                Mainloop Builder Setup
  ///////////////////////////////////////////////////////////
  
  ///////////////////////////////////////////
  // 1. Describe A and B tensors
  ///////////////////////////////////////////
  using ElementA       = // TBD
  constexpr int AlignA = // TBD
  using GmemLayoutA    = // TBD
  using ElementB       = // TBD
  constexpr int AlignB = // TBD
  using GmemLayoutB    = // TBD

  // Mma's accumulator type
  using ElementAccumulator = float;           // Always float for block scaled tcgen05.mma instructions

  //////////////////////////////////////////
  // 2. Choose Performance Parameters
  //////////////////////////////////////////

  // Tile and cluster shapes
  // Collective MMA takes tile shape of the MMA operation as input
  using KernelMainloopPolicy     = // TBD
  using MmaTileShape_MNK         = // TBD
  using ClusterShape_MNK         = // TBD

  using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
      cutlass::arch::Sm100, cutlass::arch::OpClassBlockScaledTensorOp,      // Arch and Tensorop spec
      ElementA, GmemLayoutA, AlignA,                                        // A tensor elem type, layout and alignment requirement
      ElementB, GmemLayoutB, AlignB,                                        // B tensor elem type, layout and alignment requirement
      ElementAccumulator,                                                   // Mma instruction accumulator type
      MmaTileShape_MNK, ClusterShape_MNK,                                   // Mma instruction tile shape, cluster shape
      // Epilogue's SMEM usage that needs to be subtracted from overall SMEM capacity 
      cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(sizeof(typename CollectiveEpilogue::SharedStorage))>,
      KernelMainloopPolicy                                                  // Kernel schedule policy.
                                                                            // Auto or using targeted scheduling policy
    >::CollectiveOp;

```

--------------------------------

### CUDA Kernel Function Definition

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/mma_docs/tcgen05_programming.html

Defines a CUDA kernel using the @cute.kernel decorator. It outlines the setup, partitioning, main loop for MMA operations, and epilogue for copying results back to global memory. Requires setup of TMA, SMEM, and TMEM allocators.

```python
@cute.kernel
def kernel(
    tiled_mma: cute.TiledMma,
    tma_a: cpasync.TmaInfo,
    tma_b: cpasync.TmaInfo,
    mC: cute.Tensor,
):
    # -- Setup --
    bidx, bidy, _ = cute.arch.block_idx()
    mma_coord_mnk = (bidx, bidy, None)

    # Global tensors for A and B live inside the TMA descriptor
    mA = tma_a.tma_tensor   # (M, K)
    mB = tma_b.tma_tensor   # (N, K)

    # Allocate SMEM for A, B (staged) and pipeline barriers
    smem = cutlass.utils.SmemAllocator()
    sA = smem.allocate_tensor(...)  # staged SMEM for A
    sB = smem.allocate_tensor(...)  # staged SMEM for B

    # Allocate TMEM for the accumulator
    tmem = cutlass.utils.TmemAllocator(...)
    tmem.allocate(num_cols=512)

    # -- Partition and make fragments --
    # (BM, BK, k)
    gA = cute.local_tile(mA, mma_tiler_mnk, mma_coord_mnk, proj=(1, None, 1))
    # (BN, BK, k)
    gB = cute.local_tile(mB, mma_tiler_mnk, mma_coord_mnk, proj=(None, 1, 1))
    # (BM, BN)
    gC = cute.local_tile(mC, mma_tiler_mnk, mma_coord_mnk, proj=(1, 1, None))

    thr_mma = tiled_mma.get_slice(0)
    tCgA = thr_mma.partition_A(gA)   # (MMA, MMA_M, MMA_K, num_k_tiles)
    tCgB = thr_mma.partition_B(gB)   # (MMA, MMA_N, MMA_K, num_k_tiles)
    tCgC = thr_mma.partition_C(gC)   # (MMA, MMA_M, MMA_N)

    # SMEM descriptor fragments
    tCrA = tiled_mma.make_fragment_A(sA)   # (MMA, MMA_M, MMA_K, STAGE)
    tCrB = tiled_mma.make_fragment_B(sB)   # (MMA, MMA_N, MMA_K, STAGE)

    # TMEM accumulator
    acc_shape = tiled_mma.partition_shape_C(mma_tiler_mnk[:2])
    tCtAcc = tiled_mma.make_fragment_C(acc_shape)  # (MMA, MMA_M, MMA_N)

    # Bind accumulator to TMEM
    tmem.wait_for_alloc()
    tmem_ptr = tmem.retrieve_ptr(cutlass.Float32)
    tCtAcc = cute.make_tensor(tmem_ptr, tCtAcc.layout)

    # TMA partition for global → shared memory copies
    tAsA, tAgA = cute.nvgpu.cpasync.tma_partition(
        tma_a.atom, 0, cute.make_layout(1),
        cute.group_modes(sA, 0, 3), cute.group_modes(tCgA, 0, 3),
    )
    tBsB, tBgB = cute.nvgpu.cpasync.tma_partition(
        tma_b.atom, 0, cute.make_layout(1),
        cute.group_modes(sB, 0, 3), cute.group_modes(tCgB, 0, 3),
    )

    # -- Main loop: iterate over K-tiles --
    num_k_tiles = cute.size(gA, mode=[2])
    for k_tile_idx in cutlass.range(num_k_tiles):
        # TMA load A, B into staged SMEM (producer side)
        # copy(tma_a.atom, tAgA[k], tAsA[stage])
        # copy(tma_b.atom, tBgB[k], tBsB[stage])
        # ... (see pipeline documentation)

        # Wait for data
        ab_full = ab_consumer.wait_and_advance()

        # MMA
        tiled_mma.set(tcgen05.Field.ACCUMULATE, k_tile_idx != 0)
        tile_crd = (None, None, None, ab_full.index)
        cute.gemm(tiled_mma, tCtAcc, tCrA[tile_crd], tCrB[tile_crd], tCtAcc)

        ab_full.release()

    # -- Epilogue: copy accumulator from TMEM to global memory --
    copy_atom_t2r = cute.make_copy_atom(
        tcgen05.Ld32x32bOp(tcgen05.Repetition.x64), cutlass.Float32,
    )
    tiled_copy_t2r = tcgen05.make_tmem_copy(copy_atom_t2r, tCtAcc[(None, None), 0, 0])
    thr_copy_t2r = tiled_copy_t2r.get_slice(tidx)

    tTR_tAcc = thr_copy_t2r.partition_S(tCtAcc)
    tTR_gC   = thr_copy_t2r.partition_D(tCgC)
    tTR_rAcc = cute.make_rmem_tensor(tTR_gC[None, None, 0].shape, acc_dtype)

    # TMEM → RMEM, then RMEM → GMEM
    for i in cutlass.range(num_tiles):
        cute.copy(tiled_copy_t2r, tTR_tAcc[None, None, i], tTR_rAcc)
        cute.copy(store_atom, tTR_rAcc, tTR_gC[None, None, i])

```

--------------------------------

### Initialize Half-Precision Tensor with Random Gaussian Distribution

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/utilities.html

Example demonstrating the initialization of a half-precision tensor using a random Gaussian distribution. Supports both host and device memory initialization.

```cpp
#include <cutlass/numeric_types.h>
#include <cutlass/layout/matrix.h>
#include <cutlass/util/reference/host/tensor_fill.h>
#include <cutlass/util/reference/device/tensor_fill.h>
#include <cutlass/util/host_tensor.h>

int main() {
  int rows = 128;
  int columns = 64;

  double mean = 0.5;
  double stddev = 2.0;
  uint64_t seed = 0x2019;

  // Allocate a column-major tensor with half-precision elements
  cutlass::HostTensor<cutlass::half_t, cutlass::layout::ColumnMajor> tensor({rows, columns});

  // Initialize in host memory
  cutlass::reference::host::TensorFillRandomGaussian(
    tensor.host_view(),
    seed,
    mean,
    stddev);

  // Initialize in device memory
  cutlass::reference::device::TensorFillRandomGaussian(
    tensor.device_view(),
    seed,
    mean,
    stddev);

  return 0;
}

```

--------------------------------

### cutlass.cute.front

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Recursively gets the first element of an input structure.

```APIDOC
## cutlass.cute.front

### Description
Recursively get the first element of input. This function traverses a hierarchical structure (like a layout or tensor) and returns the first element at the deepest level. It’s particularly useful for accessing the first stride value in a layout to determine properties like majorness.

### Parameters
* **input** (Union[Tensor, Layout, Stride]) – The hierarchical structure to traverse
* **loc** (source location, optional) – Source location where it’s called, defaults to None
* **ip** (insertion pointer, optional) – Insertion pointer for IR generation, defaults to None

### Returns
The first element at the deepest level of the input structure

### Return type
Union[int, float, bool, ir.Value]
```

--------------------------------

### Pipeline Consumer Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/pipeline.html

Demonstrates the usage of `PipelineConsumer` methods like `try_wait` and `wait_and_advance` for managing asynchronous data flow. It shows how to attempt a non-blocking wait and then proceed with a blocking wait if necessary, followed by releasing the buffer.

```python
pipeline = PipelineAsync.create(...)
producer, consumer = pipeline.make_participants()
for i in range(iterations):
    # Try to wait for buffer to be full
    try_wait_token = consumer.try_wait()

    # Do something else independently
    ...

    # Wait for buffer to be full & Move index to next stage
    # If try_wait_token is True, return immediately
    # If try_wait_token is False, block until buffer is full
    handle = consumer.wait_and_advance(try_wait_token)

    # Consume data
    handle.release(  )  # Signal buffer is empty

    # Alternative way to do this is:
    # handle.release()  # Signal buffer is empty

```

--------------------------------

### cutlass.utils.sm100.get_layoutSFB_TV

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils_sm100.html

Get the Thread-Value layout for scale factor B.

```APIDOC
## cutlass.utils.sm100.get_layoutSFB_TV

### Description
Get the Thread-Value layout for scale factor B.

### Signature
```python
cutlass.utils.sm100.get_layoutSFB_TV( _tiled_mma : TiledMma_, ) → cutlass.cute.typing.Layout_
```
```

--------------------------------

### cutlass.utils.sm100.get_layoutSFA_TV

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils_sm100.html

Get the Thread-Value layout for scale factor A.

```APIDOC
## cutlass.utils.sm100.get_layoutSFA_TV

### Description
Get the Thread-Value layout for scale factor A.

### Signature
```python
cutlass.utils.sm100.get_layoutSFA_TV( _tiled_mma : TiledMma_, ) → cutlass.cute.typing.Layout_
```
```

--------------------------------

### Consumer Get Barrier Pointer

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/pipeline.html

Retrieves the barrier pointer for the consumer.

```python
pipeline.consumer_get_barrier(
    state,
    _loc = loc,
    _ip = ip
)
```

--------------------------------

### Compile Kernel with Fake Tensors using TVM FFI

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/compile_with_tvm_ffi.html

This example demonstrates how to define a kernel, compile it with TVM FFI using fake tensors for shape and type information, and then execute it with actual PyTorch tensors. Ensure the compilation command includes the "--enable-tvm-ffi" option.

```python
import cutlass.cute as cute
import torch

@cute.kernel
def device_add_one(a: cute.Tensor, b: cute.Tensor):
   threads_per_block = 128
   cta_x_, _, _ = cute.arch.block_idx()
   tid_x, _, _ = cute.arch.thread_idx()
   tid = cta_x_ * threads_per_block + tid_x
   if tid < a.shape[0]:
      b[tid] = a[tid] + 1.0

@cute.jit
def add_one(a: cute.Tensor, b: cute.Tensor):
   n = a.shape[0]
   threads_per_block = 128
   blocks = (n + threads_per_block - 1) // threads_per_block
   device_add_one(a, b).launch(
      grid=(blocks, 1, 1),
      block=(threads_per_block, 1, 1),
   )

def example_add_one():
   n = cute.sym_int()
   a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   # compile the kernel with "--enable-tvm-ffi" option and example input tensors
   compiled_add_one = cute.compile(add_one, a_cute, b_cute, options="--enable-tvm-ffi")
   # now compiled_add_one is a TVM-FFI function that can be called with torch.Tensor as input
   a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
   b_torch = torch.empty(10, dtype=torch.float32, device="cuda")
   compiled_add_one(a_torch, b_torch)
   print("result of b_torch after compiled_add_one(a_torch, b_torch)")
   print(b_torch)

```

--------------------------------

### Producer Get Barrier Pointer

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/pipeline.html

Retrieves the barrier pointer for the producer.

```python
pipeline.producer_get_barrier(
    state,
    _loc = loc,
    _ip = ip
)
```

--------------------------------

### Compile Kernel with Tuple Arguments using TVM FFI

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/compile_with_tvm_ffi.html

Demonstrates compiling a kernel that accepts a tuple containing tensors and a float, using the `--enable-tvm-ffi` option. This example shows how to define a kernel that takes a tuple as input and how to compile and run it.

```python
import torch
from cutlass import cute
from typing import Tuple

@cute.kernel
def device_add_one(a: cute.Tensor, b: cute.Tensor, c: cute.Float32):
   threads_per_block = 128
   cta_x_, _, _ = cute.arch.block_idx()
   tid_x, _, _ = cute.arch.thread_idx()
   tid = cta_x_ * threads_per_block + tid_x
   if tid < a.shape[0]:
      b[tid] = a[tid] + c

@cute.jit
def add_one_with_tuple(a: Tuple[cute.Tensor, cute.Tensor, cute.Float32]):
   n = a[0].shape[0]
   threads_per_block = 128
   blocks = (n + threads_per_block - 1) // threads_per_block
   device_add_one(a[0], a[1], a[2]).launch(grid=(blocks, 1, 1), block=(threads_per_block, 1, 1))

def example_add_one_with_tuple():
   n = cute.sym_int()
   a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   compiled_add_one = cute.compile(
      add_one_with_tuple, (a_cute, b_cute, cute.Float32(4)),
      options="--enable-tvm-ffi"
   )
   a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
   b_torch = torch.empty(10, dtype=torch.float32, device="cuda")
   compiled_add_one((a_torch, b_torch, 5))
   print("result of b_torch after compiled_add_one((a_torch, b_torch, 5))")
   print(b_torch)

example_add_one_with_tuple()

```

--------------------------------

### Get Device Multiprocessor Count

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils.html

Retrieves the number of multiprocessors on the CUDA device.

```python
def get_device_multiprocessor_count() -> int:
    pass
```

--------------------------------

### Get L2 Cache Size

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils.html

Retrieves the L2 cache size in bytes.

```python
def get_l2_cache_size_in_bytes() -> int:
    pass
```

--------------------------------

### Get Initial Work Tile Info

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils.html

Retrieves the initial work tile information.

```APIDOC
## initial_work_tile_info

### Description
Retrieves the initial work tile information.

### Parameters
* **loc** (_cutlass._mlir.ir.Location | None_) – Optional location for MLIR operations.
* **ip** (_cutlass._mlir.ir.InsertionPoint | None_) – Optional insertion point for MLIR operations.

### Returns
A WorkTileInfo object representing the initial work.

### Return type
WorkTileInfo
```

--------------------------------

### CopyBulkG2SMulticastOp Initialization

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute_nvgpu_cpasync.html

Initializes a CopyBulkG2SMulticastOp for bulk asynchronous GMEM to SMEM multicast copy operations.

```python
cutlass.cute.nvgpu.cpasync.CopyBulkG2SMulticastOp()
```

--------------------------------

### Get Current Work Tile Info

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils.html

Retrieves the current work tile information.

```APIDOC
## get_current_work

### Description
Retrieves the current work tile information.

### Parameters
* **loc** (_cutlass._mlir.ir.Location | None_) – Optional location for MLIR operations.
* **ip** (_cutlass._mlir.ir.InsertionPoint | None_) – Optional insertion point for MLIR operations.

### Returns
A WorkTileInfo object representing the current work.

### Return type
WorkTileInfo
```

--------------------------------

### Row-Major Layout Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/layout.html

Illustrates the CUTLASS `layout::RowMajor` for representing a row-major matrix layout.

```cpp
layout::RowMajor layout(lda); 

int offset = layout({row, column});     // returns lda * row + column

```

--------------------------------

### Build CUTLASS with Heuristics Options

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/heuristics.html

Configure and build CUTLASS using CMake, specifying paths to the problem list JSON and the desired number of configurations per problem. Hardware is auto-detected.

```bash
$ cmake .. \
    -DCUTLASS_NVCC_ARCHS=90a \
    -DCUTLASS_LIBRARY_HEURISTICS_PROBLEMS_FILE=<path_to_your_problem_list.json> \
    -DCUTLASS_LIBRARY_HEURISTICS_CONFIGS_PER_PROBLEM=<number of configurations to build per problem> 

$ make cutlass_profiler -j

```

--------------------------------

### cutlass.utils.make_smem_layout_a

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/utils.html

Helps in getting the partitioned shape of the A tensor and selecting the heuristic SMEM layout atom.

```APIDOC
## cutlass.utils.make_smem_layout_a

### Description
This function helps in getting the partitioned shape of the A tensor based on the tiled_mma & MMA tiler, selecting the heuristic SMEM layout atom, tiling it to the MMA tile shape, and staging the SMEM layout.

### Parameters
* **tiled_mma** (_TiledMma_) – The tiled MMA operation.
* **mma_tiler_mnk** (_cute.Tile_) – The MMA tiler for MNK dimensions.
* **a_dtype** (_Type_ _[__Numeric_ _]_) – Element type of the A tensor.
* **num_stages** (_int_) – Number of pipeline stages.
* **is_k_major** (_bool | None_) – Flag indicating if K is the major dimension.

### Returns
Staged SMEM layout for the A tensor.

### Return type
Union[cute.Layout, cute.ComposedLayout]
```

--------------------------------

### Example: Add One with Explicit Stream

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/compile_with_tvm_ffi.html

Demonstrates compiling and executing a kernel with an explicit CUDA stream. It uses fake tensors and streams for compilation, then real PyTorch tensors and streams for execution.

```python
def example_add_one_with_stream():
   n = cute.sym_int()
   a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
   # Fake stream is a placeholder for stream argument
   stream = cute.runtime.make_fake_stream()
   compiled_add_one = cute.compile(
      add_one_with_stream, a_cute, b_cute, stream, options="--enable-tvm-ffi"
   )
   a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
   b_torch = torch.empty(10, dtype=torch.float32, device="cuda")
   torch_stream = torch.cuda.current_stream()
   compiled_add_one(a_torch, b_torch, torch_stream)
   torch_stream.synchronize()
   print("result of b_torch after compiled_add_one(a_torch, b_torch, torch_stream)")
   print(b_torch)
```

--------------------------------

### make_tiled_copy

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Create a tiled type given a TV partitioner and tiler.

```APIDOC
## make_tiled_copy(_atom: Any, _layout_tv: Any, _tiler_mn: Any, *_, _loc: Location | None = None, _ip: InsertionPoint | None = None) -> TiledCopy

### Description
Create a tiled type given a TV partitioner and tiler.

### Parameters

#### Path Parameters
* **atom** (Any) - The atom to use for creating the tiled copy.
* **layout_tv** (Any) - The TV partitioner layout.
* **tiler_mn** (Any) - The MN tiler.

#### Optional Parameters
* **loc** (Location | None) - Source location for MLIR, defaults to None.
* **ip** (InsertionPoint | None) - Insertion point, defaults to None.

### Returns
A tiled copy.

### Return type
TiledCopy
```

--------------------------------

### Access Union Size and Alignment

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_api/cute.html

Statically get the size and alignment of a C union defined with the cute.union decorator.

```python
size = data_union.__sizeof__()
align = data_union.__alignof__()
```