NVIDIA CUTLASS

https://docs.nvidia.com/cutlass/
Admin
NVIDIA CUTLASS is a collection of CUDA C++ template abstractions and Python domain-specific...
Tokens:390,663
Snippets:3,681
Trust Score:10
Update:2 months ago
Show doc for...
Tokens:
Raw
### Install Recommended Development Dependencies

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

Install common tools like PyTorch and Jupyter for development and running examples.

```bash
pip install torch jupyter mypy==1.19.1
```

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

### Display Convolution Example Help

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

Command to view the command-line interface options for the convolution example.

```bash
$ ./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop --help
```

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

### Clone and Setup CUTLASS DSL from Repository

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

Use the setup script from the repository to ensure compatibility with specific CUDA toolkit versions.

```bash
git clone https://github.com/NVIDIA/cutlass.git

# For CUDA Toolkit 12.9:
./cutlass/python/CuTeDSL/setup.sh --cu12

# For CUDA Toolkit 13.3:
./cutlass/python/CuTeDSL/setup.sh --cu13
```

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

### Configure and Build CUTLASS Convolution Example

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

Commands to create a build directory, configure the project for compute capability 7.5, and compile the specific convolution example.

```bash
$ mkdir build && cd build

$ cmake .. -DCUTLASS_NVCC_ARCHS=75
```

```bash
$ make 09_turing_tensorop_conv2dfprop

$ ls examples/09_turing_tensorop_conv2dfprop 
examples/09_turing_tensorop_conv2dfprop
```

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

### Install Preview CUTLASS DSL via Pip

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

Install preview releases from the NVIDIA PyPI index for specific CUDA toolkits.

```bash
# For CUDA Toolkit 12.9:
pip install --pre nvidia-cutlass-dsl --extra-index-url https://pypi.nvidia.com

# For CUDA Toolkit 13.3:
pip install --pre "nvidia-cutlass-dsl[cu13]" --extra-index-url https://pypi.nvidia.com
```

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

### Execute Convolution Example

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

Commands to run the convolution example with specific tensor dimensions and optional reference checking.

```bash
$ ./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop  --n=32 --h=224 --w=224 --c=128 --k=256 --r=1 --s=1

$ ./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop  --n=1 --h=224 --w=224 --c=32 --k=32 --r=3 --s=3 --ref-check
```

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

### Build and Run with synclog

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

Build the target example and execute it with iterations set to 0 to capture synchronization logs.

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

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

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

### Sample synclog Output

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

Example of the text format produced in the synclog output file.

```text
synclog start
synclog at 1: cluster_barrier_init line=281 time=1725400116233388736 thread=0,0,0 block=0,0,0 smem_addr=197632 arrive_count=1
synclog at 13: fence_barrier_init line=583 time=1725400116233388768 thread=32,0,0 block=0,0,0 
...
```

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

### Verify IKET Installation

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

Check if the run-iket tool is available in the current environment.

```bash
run-iket --help
```

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

### Create and manipulate Layouts

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

Examples of creating layouts with specific shapes and strides, and performing coordinate-to-index mapping.

```python
# Creating a layout with shape (4,8) and default stride (column major)
layout = cute.make_layout((4, 8))

# Creating a layout with explicit shape and stride (row major)
layout = cute.make_layout((4, 8), stride=(8, 1))

# Accessing layout properties
shape = layout.shape      # Returns (4, 8)
stride = layout.stride    # Returns (8, 1)

# Mapping a coordinate to an index: (2, 3) -> 2 * 8 + 3 * 1 = 19
idx = cute.crd2idx((2, 3), layout)
```

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

### Benchmark Output Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/tutorials/003_host_latency_best_practices.html

Sample output showing the performance improvement achieved using CUDA Graphs.

```text
[20 launches without CUDA Graph] avg of 1 iterations: 7.496e-04 seconds
[20 launches with CUDA Graph   ] avg of 1 iterations: 1.972e-04 seconds
Speedup with CUDA Graph: 3.80x
```

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

### iket.range_start(name, payload=None)

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

Starts a token-based range and returns a token used to close the range.

```APIDOC
## iket.range_start(name, payload=None)

### Description
Starts a token-based range and returns a token. Closed by iket.range_end(token).

### Parameters
- **name** (string) - Required - The name of the range.
- **payload** (numeric) - Optional - A numeric value attached to the start event.
```

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

### Example output of operator discovery

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/tutorials/000_gemm.html

Shows the console output resulting from the operator discovery process.

```text
A total of 184434 Operator instances are available.
Of these, 572 support the given arguments.
Of these, 548 support the given GPU architecture.
Picked operator with name: cutedsl.PersistentDenseGemmOperator_sm100_ttt_AFloat16_BFloat16_outFloat16_accFloat32_2cta_cluster2x1x1_tile128x32x64_tma_store
```

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

### Local clangd Configuration

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

Example local configuration for per-project settings, requiring absolute paths for include directories.

```yaml
CompileFlags:
  Add:
    - -I</absolute/path/to/cutlass>/include/
    - -I</absolute/path/to/cutlass>/tools/util/include/
    - -I</absolute/path/to/cutlass>/cutlass/examples/common/
```

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

### Control Flow Loop Examples

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

Demonstrates the usage of range, cutlass.range, and cutlass.range_constexpr for different loop evaluation strategies.

```python
@cute.jit
def control_flow_examples(bound: cutlass.Int32):
    n = 10

    # ✅ This loop is Python loop, evaluated at compile time.
    for i in cutlass.range_constexpr(n):
        cute.printf("%d\n", i)

    # ✅ This loop is dynamic, even when bound is Python value.
    for i in range(n):
        cute.printf("%d\n", i)

    # ❌ This loop bound is a dynamic value, not allowed in Python loop.
    # Should use `range` instead.
    for i in cutlass.range_constexpr(bound):
        cute.printf("%d\n", i)

    # ✅ This loop is dynamic, emitted IR loop.
    for i in range(bound):
        cute.printf("%d\n", i)

    # ✅ This loop is dynamic, emitted IR loop with unrolling
    for i in cutlass.range(bound, unroll=2):
        cute.printf("%d\n", i)
```

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

### Example output for static layout (3):(1)

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

The expected output when running the compiled function with a 3-element tensor.

```text
  tensor.layout: (3):(1)
  tensor: raw_ptr(0x00000000079e5100: i16, generic, align<2>) o (3):(1) =
( 1, 2, 3 )
```

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

### Install Heuristics Dependency

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

Install the required Python package for heuristic calculations.

```bash
pip install nvidia-matmul-heuristics
```

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

### Install CUTLASS Operator variants

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/overview.html

Install the core package or the full development suite.

```bash
# Install only nvidia-cutlass-operators core
pip install nvidia-cutlass-operators

# Install all dependencies to develop, run tests, etc.
pip install nvidia-cutlass-operators[dev]
```

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

### setup_resources_and_tasks()

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_api/task_manager.html

Initializes pipelines and barriers for every resource, including SMEM allocation if configured.

```APIDOC
## setup_resources_and_tasks()

### Description
Materializes and initializes pipelines and barriers for every resource. This method must be called exactly once before `run()` to allocate SMEM and instantiate pipeline objects.

### Returns
- **None**
```

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

### Install CUTLASS Operator with PyTorch

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/overview.html

Use this command to install the package with PyTorch support.

```bash
pip install nvidia-cutlass-operators[torch]
```

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

### TaskManager.setup_resources_and_tasks

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_api/task_manager.html

Materializes and initializes pipelines and barriers for every resource.

```APIDOC
## TaskManager.setup_resources_and_tasks()

### Description
Materializes and initializes pipelines and barriers for every resource by calling resource.create(). This must be called once before run().
```

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

### Install TVM FFI dependencies

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

Commands to install the required TVM FFI package and optional performance extensions.

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

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

### Initialize environment and check device support

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/tutorials/004_fake_tensors.html

Imports necessary modules and verifies that the current environment supports the required GPU compute capability.

```python
import torch

import cutlass.operators as ops

torch.manual_seed(2025)

if not (status := ops.utils.device.device_or_env_supports("80")):
    print(
        f"This notebook requires a GPU with compute capability >= 80.\n{status.error}"
    )
    import sys

    sys.exit(0)
```

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

### Print Tensor Output Example

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

Example output format for the print_tensor function, showing metadata and data values.

```text
tensor(raw_ptr<@..., Float32, generic, align(4)> o (8,5):(5,1), data=
       [[-0.4326, -0.5434,  0.1238,  0.7132,  0.8042],
        [-0.8462,  0.9871,  0.4389,  0.7298,  0.6948],
        [ 0.3426,  0.5856,  0.1541,  0.2923,  0.6976],
        [-0.1649,  0.8811,  0.1788,  0.1404,  0.2568],
        [-0.2944,  0.8593,  0.4171,  0.8998,  0.1766],
        [ 0.8814,  0.7919,  0.7390,  0.4566,  0.1576],
        [ 0.9159,  0.7577,  0.6918,  0.0754,  0.0591],
        [ 0.6551,  0.1626,  0.1189,  0.0292,  0.8655]])
```

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

### TaskManager.setup_resources_and_tasks()

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_api/task_manager.html

Configures the resources and tasks required for execution.

```APIDOC
## TaskManager.setup_resources_and_tasks()

### Description
Prepares the environment by setting up necessary resources and tasks. This must be called before other operations.

### Raises
- **AssertionError** - If called incorrectly or if resources fail to initialize.
```

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

### Print Owning Tensor Output

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

Example output from calling print on the owning tensors created in the previous example.

```text
rmem_4x8_col  : ptr[32b](0x7fff48929460) o (_4,_8):(_1,_4)
rmem_4x8_row  : ptr[32b](0x7fff489294e0) o (_4,_8):(_8,_1)
rmem_4x8_pad  : ptr[32b](0x7fff489295e0) o (_4,_8):(_32,_2)
rmem_4x8_like : ptr[32b](0x7fff48929560) o (_4,_8):(_8,_1)
```

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

### Print Nonowning Tensor Output

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

Example output from calling print on the nonowning tensors created in the previous example.

```text
tensor_8     : ptr[32b](0x7f42efc00000) o _8:_1
tensor_8s    : ptr[32b](0x7f42efc00000) o _8:_1
tensor_8d2   : ptr[32b](0x7f42efc00000) o 8:2
gmem_8s      : gmem_ptr[32b](0x7f42efc00000) o _8:_1
gmem_8d      : gmem_ptr[32b](0x7f42efc00000) o 8:_1
gmem_8sx16d  : gmem_ptr[32b](0x7f42efc00000) o (_8,16):(_1,_8)
gmem_8dx16s  : gmem_ptr[32b](0x7f42efc00000) o (8,_16):(_16,_1)
smem_4x8_col : smem_ptr[32b](0x7f4316000000) o (_4,_8):(_1,_4)
smem_4x8_row : smem_ptr[32b](0x7f4316000000) o (_4,_8):(_8,_1)
```

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

### Initialize and Use PipelineAsync

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

Demonstrates creating a pipeline with a specified number of stages and performing producer-consumer operations using acquire, commit, wait, and release methods.

```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
```

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

### Enable Unity Build

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

Reduces binary size and avoids linker limitations by instantiating multiple kernel instances in each compilation unit.

```bash
$ cmake .. -DCUTLASS_NVCC_ARCHS="70;75;80" -DCUTLASS_LIBRARY_KERNELS=all -DCUTLASS_UNITY_BUILD_ENABLED=ON
```

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

### Install Stable CUTLASS DSL via Pip

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

Install the stable release of the CUTLASS DSL package for specific CUDA toolkits.

```bash
# For CUDA Toolkit 12.9:
pip install nvidia-cutlass-dsl

# For CUDA Toolkit 13.3:
pip install "nvidia-cutlass-dsl[cu13]"
```

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

### Initialize benchmarking data and operator

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/tutorials/003_host_latency_best_practices.html

Sets up tensors, GEMM arguments, and selects an operator for benchmarking.

```python
warmup_iterations = 10
profiling_iterations = 100
total_iterations = warmup_iterations + profiling_iterations

# Use a small problem size to showcase host overheads
L, M, N, K = 1, 128, 128, 128

# We use different operands in each iteration. Though not particularly relevant for
# host latency, this is a best practice when benchmarking GPU kernels to avoid
# unrealistic caching effects.
As = [
    torch.randint(-1, 2, (M, K), device="cuda", dtype=torch.float16)
    for _ in range(total_iterations)
]
Bs = [
    torch.randint(-1, 2, (K, N), device="cuda", dtype=torch.float16)
    for _ in range(total_iterations)
]
outs = [
    torch.empty((M, N), device="cuda", dtype=torch.float16)
    for _ in range(total_iterations)
]

# Construct arguments outside of the benchmarking loop. We will later also consider
# cases in which they are constructed inside the benchmarking loop.
args = [
    ops.GemmArguments(A=As[i], B=Bs[i], out=outs[i], accumulator_type=torch.float32)
    for i in range(total_iterations)
]

references = [(As[i] @ Bs[i]).to(outs[i].dtype) for i in range(total_iterations)]

target_sm = ops.utils.device.device_or_env_target_sm()
operators = ops.get_operators(args[0], target_sm=target_sm)
assert len(operators) > 0

operator = operators[0]
```

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

### Implementing a 4-stage Asynchronous Pipeline

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

Demonstrates the setup and usage of a PipelineAsync instance with producer and consumer thread roles.

```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;
  }
}
```

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

### Define SMEM Layout Parameters for F16 Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/guides/mma/wmma_programming.html

Example parameters for A and B operands in an F16 GEMM kernel with specific tiling and swizzle configurations.

```text
A operand (row-major, smem_tiler = (128, 32, 4)):
  major_mode_size = 32
  atom = (8, 32):(32, 1)
  swizzle = make_swizzle(2, 4, 3)
  tiled layout -> sA: (128, 32, 4)

B operand (col-major, smem_tiler = (128, 32, 4)):
  major_mode_size = min(128, 64) = 64
  atom = (64, 8):(1, 64)
  swizzle = make_swizzle(3, 4, 3)
  tiled layout -> sB: (128, 32, 4)
```

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

### Configure Profiler for Full Kernel Instantiation

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

CMake configuration to enable all kernel sizes and architectures, followed by the build command.

```bash
$ cmake .. -DCUTLASS_NVCC_ARCHS="70;75;80" -DCUTLASS_LIBRARY_KERNELS=all  -DCUTLASS_UNITY_BUILD_ENABLED=ON
...
$ make cutlass_profiler -j
```

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

### Configure Build Directory and Architectures

Source: https://docs.nvidia.com/cutlass/latest/overview.html

Create a build directory and run CMake with specific architecture flags to limit compilation time.

```bash
$ mkdir build && cd build

$ cmake .. -DCUTLASS_NVCC_ARCHS=80               # compiles for NVIDIA's Ampere Architecture
```

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

### Output of ArithmeticTupleIterator

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

The resulting output of the ArithmeticTupleIterator example.

```text
(42,7,_9)
```

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

### Allocate HostTensor

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

Example of allocating a column-major matrix with single-precision elements.

```cpp
#include <cutlass/layout/matrix.h>
#include <cutlass/util/host_tensor.h>

int main() {
  int rows = 32;
  int columns = 16;

  cutlass::HostTensor<float, cutlass::layout::ColumnMajor> tensor({rows, columns});

  return 0;
}
```

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

### TS Kernel Authoring Skeleton

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_general/ts_introduction.html

A comprehensive example demonstrating the declaration of resource classes, dependency graph definition, schedule capture, task binding, and validation.

```python
# 1. Resource classes declare the schedule-facing operations.
class GmemAResource(MemoryResource):
    @consumer_work
    def compute_coords(self, stage_info):
        ...

class SmemAResource(MemoryResource):
    @producer_work
    def tma_load_a(self, stage_info, coord_k, coord_m):
        ...

# 2. The kernel instantiates resources, optionally with pipeline configs.
gmem_a = GmemAResource(name="gmem_a")
smem_a = SmemAResource(
    name="smem_a",
    pipeline_config=smem_a_pipeline_config,
)

# 3. The dependency graph states the intended resource ordering.
resource_dependency_graph = {
    smem_a: [gmem_a],
}

# 4. The captured schedule records the task ordering.
@schedule
def load_a_schedule(gmem_a: GmemAResource,
                    smem_a: SmemAResource) -> None:
    with domain_loop(0, num_k_tiles, 1):
        coord_k, coord_m, coord_n = gmem_a.compute_coords()
        smem_a.acquire()
        smem_a.tma_load_a(coord_k=coord_k, coord_m=coord_m)
        smem_a.commit()

# 5. The task binds resources, schedule, and warp assignment.
load_a_task = Task(
    name="load_a",
    src_resources=[gmem_a],
    dst_resources=[smem_a],
    warp_idx=0,
    num_warps=1,
    schedule=load_a_schedule(gmem_a, smem_a),
)

# 6. The manager validates the schedule before lowering.
task_manager = TaskManager(
    tasks=[load_a_task],
    resource_dependency_graph=resource_dependency_graph,
)
task_manager.print_and_verify()

# 7. Fill the @consumer_work / @producer_work bodies, then run
#    the kernel correctness check against the expected result.
```

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

### Create and use ScaledBasis elements

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

Demonstrates initializing ScaledBasis with various scales and modes, and applying them within a layout to map coordinates to indices.

```python
# Create a scaled basis with integer scale and mode
sb1 = ScaledBasis(2, 0)  # 2 * E(0)

# Create a scaled basis with a Ratio scale
sb2 = ScaledBasis(Ratio(1, 2), 1)  # (1/2) * E(1)

# Create a scaled basis with a list of modes
sb3 = ScaledBasis(4, [0, 1])  # 4 * E([0, 1])

# Scaled basis elements are commonly used in layout strides
layout = make_layout((4, 8), stride=(ScaledBasis(2, 0), ScaledBasis(1, 1)))

# This creates a layout with strides (2@0, 1@1) representing
# a coordinate system where each dimension has its own basis

# Example: Mapping coordinates to indices using the layout
coord = (2, 3)
idx = crd2idx(coord, layout)  # Maps (2, 3) to (4, 3)
```

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

### Displaying a CuTe Tensor

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

Example of a printed CuTe tensor representation.

```text
ArithTuple(0,_0,_0,_0) o ((_128,_64),2,3,1):((_1@0,_1@1),_64@1,_1@2,_1@3)
```

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

### Constructing Layouts with make_layout

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/01_layout.html

Demonstrates various ways to construct Layout objects using static and dynamic integers, shapes, and strides.

```cpp
Layout s8 = make_layout(Int<8>{});
Layout d8 = make_layout(8);

Layout s2xs4 = make_layout(make_shape(Int<2>{},Int<4>{}));
Layout s2xd4 = make_layout(make_shape(Int<2>{},4));

Layout s2xd4_a = make_layout(make_shape (Int< 2>{},4),
                             make_stride(Int<12>{},Int<1>{}));
Layout s2xd4_col = make_layout(make_shape(Int<2>{},4),
                               LayoutLeft{});
Layout s2xd4_row = make_layout(make_shape(Int<2>{},4),
                               LayoutRight{});

Layout s2xh4 = make_layout(make_shape (2,make_shape (2,2)),
                           make_stride(4,make_stride(2,1)));
Layout s2xh4_col = make_layout(shape(s2xh4),
                               LayoutLeft{});
```

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

### Create and use PipelineOrder

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

Demonstrates initializing a PipelineOrder with specific group and stage configurations and using it within a processing loop.

```python
# Create pipeline order with 3 groups, each with 2 stages
pipeline_order = PipelineOrder.create(
    barrier_storage=smem_ptr,      # shared memory pointer for barriers
    depth=2,                       # 2 stages per group
    length=3,                      # 3 groups total
    group_id=0,                    # current group ID (0, 1, or 2)
    producer_group=producer_warp   # cooperative group for producers
)

# In the pipeline loop
for stage in range(num_stages):
    pipeline_order.wait()          # Wait for previous group to complete
    # Process current stage
    pipeline_order.arrive()        # Signal completion to next group
```

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

### Visualize Accumulator Shape

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/guides/mma/wgmma_programming.html

Example output shape and memory layout for the accumulator.

```text
tCgC: (MMA, MMA_M=1, MMA_N=1)

make_rmem_tensor(tCgC.shape, Float32) -> acc: (MMA, 1, 1)

The accumulator stays in RMEM for the entire main loop.
cute.gemm() reads A/B from SMEM descriptors and accumulates into acc.

+-----------------------------------+
|  acc: (MMA, 1, 1) in RMEM         |
|  64 x 256 elements per warpgroup  |
|  Float32                          |
+-----------------------------------+
```

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

### Assemble a CUTLASS 3.x GEMM Kernel

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

Demonstrates the four-step process of building a GEMM kernel by defining the collective mainloop, epilogue, kernel composition, and device adapter.

```cpp
// Step 1: Generate the required collective layer mainloop specialization
using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
    ArchTag, OperatorClass,
    ElementA, LayoutA, AlignmentA,
    ElementB, LayoutB, AlignmentB,
    ElementAccumulator,
    TilesShape, ClusterShape,
    cutlass::gemm::collective::StageCountAuto,
    cutlass::gemm::collective::KernelScheduleAuto
  >::CollectiveOp;

// Step 2: Specify the collective layer epilogue type
using CollectiveEpilogue = cutlass::epilogue::collective::DefaultEpilogue<
    ElementC,
    cutlass::gemm::TagToStrideC_t<LayoutC>,
    cutlass::gemm::TagToStrideC_t<LayoutC>,
    cutlass::epilogue::thread::LinearCombination<ElementC, 1, ElementAccumulator, ElementAccumulator>>;

// Step 3: Compose the mainloop and epilogue together at the kernel layer
using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
    cute::Shape<int,int,int,int>, // ProblemShape [M,N,K,L]
    CollectiveMainloop,
    CollectiveEpilogue
>;

// Step 4: Wrap up the kernel::GemmUniversal kernel class
// with the device adapter to obtain a host-side handle to the kernel
using GemmHandle = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;
```

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

### Initialize CUTLASS environment

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/tutorials/000_gemm.html

Imports necessary modules and verifies that the current GPU meets the minimum compute capability requirements.

```python
import torch

import cutlass

import cutlass.operators as ops

if not (status := ops.utils.device.device_or_env_supports("80")):
    print(
        f"This notebook requires a GPU with compute capability >= 80.\n{status.error}"
    )
    import sys

    sys.exit(0)
```

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

### Basic Coalesce Example

Source: https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/02_layout_algebra.html

Demonstrates simplifying a layout by reducing its modes while maintaining the same mapping.

```cpp
auto layout = Layout<Shape <_2,Shape <_1,_6>>,
                     Stride<_1,Stride<_6,_2>>>{};
auto result = coalesce(layout);    // _12:_1
```

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

### SMEM Allocator Lifecycle

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_general/ts_allocators.html

Demonstrates the standard flow for initializing an SMEM allocator and integrating it with a TaskManager.

```python
allocator = SmemAllocator()
allocator.add_resource(smem_resource)
allocator.compute_layout()

task_manager = TaskManager(
    tasks=[...],
    resource_dependency_graph={...},
    smem_allocator=allocator,
)
task_manager.setup_resources_and_tasks()
```

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

### Output of implicit integer tensor

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

The resulting output of the implicit integer tensor example.

```text
counting_iter(42) o (4,5):(_1,4):
   42   46   50   54   58
   43   47   51   55   59
   44   48   52   56   60
   45   49   53   57   61
```

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

### task_scheduling.task.__init__

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_api/task.html

Initializes a new task instance with resource definitions, warp assignments, and scheduling metadata.

```APIDOC
## task_scheduling.task.__init__

### Description
Initializes a task with required resource mappings and warp assignment metadata.

### Parameters
- **src_resources** (List[MemoryResource]) - Required - Resources read by this task on the consumer side.
- **dst_resources** (List[MemoryResource]) - Required - Resources written by this task on the producer side.
- **warp_idx** (int) - Required - Index of the first warp assigned to this task.
- **num_warps** (int) - Required - Number of consecutive warps assigned to this task.
- **schedule** (object) - Required - Captured result returned by a @schedule function.
- **num_registers** (int) - Optional - Per-task register budget.
- **name** (str) - Optional - Human-readable task name.
- **debug_print** (bool) - Optional - Emit per-step debug prints.
- **run_only_on_cta_id** (int) - Optional - Hoist task selection to a specific CTA rank.
```

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

### TMA Store Interface

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

Example structure for a 3D TMA store operation in CuTe.

```cpp
struct SM90_TMA_STORE_3D {
  CUTE_DEVICE static void
  copy(void const* const desc_ptr,
       void const* const smem_ptr,
       int32_t const& crd0, int32_t const& crd1, int32_t const& crd2) {
    // ... invoke CUDA PTX instruction ...
  }
};
```

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

### @cute.union

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

A decorator used to define C-style unions where all members share the same starting offset.

```APIDOC
## @cute.union

### Description
Defines a union where all members start at offset 0. The size is the maximum size of all members, and alignment is the maximum alignment of all members.

### Methods
- `__sizeof__()`: Returns the size of the union in bytes.
- `__alignof__()`: Returns the alignment of the union in bytes.
```

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

### cutlass.cute.arch.get_min_tmem_alloc_cols

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

Get the minimum tensor memory allocation columns for a given compute capability.

```APIDOC
## cutlass.cute.arch.get_min_tmem_alloc_cols(compute_capability)

### Description
Returns the minimum TMEM allocation columns available for the specified GPU compute capability.

### Parameters
- **compute_capability** (str) - Required - The compute capability string (e.g., 'sm_100').

### Returns
- **min_cols** (int) - The minimum TMEM allocation columns.
```

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

### cutlass.cute.arch.get_max_tmem_alloc_cols

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

Get the maximum tensor memory capacity in columns for a given compute capability.

```APIDOC
## cutlass.cute.arch.get_max_tmem_alloc_cols(compute_capability)

### Description
Returns the maximum TMEM capacity in columns available for the specified GPU compute capability.

### Parameters
- **compute_capability** (str) - Required - The compute capability string (e.g., 'sm_100').

### Returns
- **capacity** (int) - The TMEM capacity in columns.
```

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

### Using Keyword Arguments and Defaults with cute.compile

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

Demonstrates defining a kernel with default arguments and invoking the compiled function using both positional and keyword arguments.

```python
import torch
from cutlass import cute

@cute.kernel
def device_add_scalar(a: cute.Tensor, b: cute.Tensor, offset: cutlass.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] + offset

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

def example_kwargs_and_defaults():
   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_constant = cute.compile(add_constant, a_cute, b_cute, 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_constant(a_torch, b_torch)
   print("result of b_torch after compiled_add_constant(a_torch, b_torch)")
   print(b_torch)
   compiled_add_constant(a_torch, b_torch, offset=4)
   print("result of b_torch after compiled_add_constant(a_torch, b_torch, offset=4)")
   print(b_torch)
```

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

### cutlass.cute.arch.get_dyn_smem_size

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

Gets the size in bytes of the dynamic shared memory specified at kernel launch.

```APIDOC
## cutlass.cute.arch.get_dyn_smem_size()

### Description
Returns the size in bytes of the dynamic shared memory allocated at kernel launch time.

### Returns
- **size** (int) - The size of dynamic shared memory in bytes.
```

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

### Access Conv2d Help

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

Display the help menu for 2D convolution operations within the profiler.

```bash
$ ./tools/profiler/cutlass_profiler --help --operation=Conv2d

Conv2d
```

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

### Install Clang dependencies on Ubuntu

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

Required packages for building with Clang on Ubuntu 22.04 LTS.

```bash
$ sudo apt-get install clang cmake ninja-build pkg-config libgtk-3-dev liblzma-dev libstdc++-12-dev
```

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

### Perform Tensor Reductions

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

Examples of reducing tensor dimensions using specified reduction profiles.

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

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

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

### Exhaustive Interleaving Checker Output

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_general/ts_printout.html

Example of the BFS exploration output generated when exhaustive_deadlock_race_check is enabled.

```text
========================================================================
Exhaustive interleaving checker — BFS exploration
========================================================================
  Task 0 (LoadTask): [ConsumerWork(InputGmem), ProducerAcquire(Smem), ...]
  Task 1 (StoreTask): [ConsumerWait(Smem), ConsumerWork(Smem), ...]
  Alias pairs:
    TmemSP <-> TmemVec
------------------------------------------------------------------------
State #7  cursors=(5, 7)  ✓ COMPLETE
  timeline (12 steps):
    t  LoadTask              StoreTask
    ─  ────────              ─────────
    0  ConsumerWork(InGmem)  ·
    1  ProducerAcquire(Smem) ·
    2  ·                     ConsumerWait(Smem)
    ...
========================================================================
BFS complete: 7 states explored, 1 complete, 0 deadlock(s), 0 race(s), 0 PDL order violation(s)
Result: SAFE
========================================================================
```

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

### cutlass.cute.make_tiled_copy(atom, layout_tv, tiler_mn)

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

Creates a tiled copy given a TV partitioner and tiler.

```APIDOC
## cutlass.cute.make_tiled_copy(atom, layout_tv, tiler_mn)

### Description
Creates a tiled copy using a specified TV partitioner and tiler.

### Parameters
- **atom** (CopyAtom) - Required - Copy atom
- **layout_tv** (Layout) - Required - Thread-value layout
- **tiler_mn** (Tiler) - Required - Tile size

### Returns
- **TiledCopy** - A tiled copy for the partitioner
```

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

### Define a resource dependency graph

Source: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/ts_api/task_manager.html

Example of mapping downstream resources to their upstream dependencies for a pipeline.

```python
resource_dependency_graph = {
    smem_ab: [gmem_ab],
    tmem_c:  [smem_ab],
    gmem_d:  [tmem_c],
}
```

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

### Convert PyTorch Tensor to CuTe

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

Example demonstrating the conversion of a PyTorch tensor to a CuTe tensor.

```python
import torch
import cutlass
from cutlass.cute.runtime import from_dlpack

x = torch.randn(30, 20, device="cpu")
y = from_dlpack(x)
```

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

### Implementing the Operator __init__ method

Source: https://docs.nvidia.com/cutlass/latest/media/docs/operators/tutorials/002_bring_your_own_kernel.html

Example of initializing the operator by extracting tile shape metadata to construct the implementation object.

```python
def __init__(self, metadata: ops.OperatorMetadata):
    # Using Python-2-style super() because we're defining this method outside of the class definition.
    super(F64GemmOperator, self).__init__(metadata)
    cta_tile_shape_mn = metadata.design.tile_shape[:2]
    self.impl = F64GemmKernelImplementation(cta_tile_shape_mn)
```

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

### Static run() API for GemmUniversalAdapter

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

Provides a static entry point for executing kernels, allowing manual management of parameters and bypassing argument lowering.

```cpp
// Primary run() entry point API that is static allowing users to create and manage their own params.
static Status
run(Params& params, cudaStream_t stream = nullptr);
```

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

### TmemBufferPool.allocate

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

Allocates a contiguous region of TMEM columns from the pool and returns a pointer to the start of that region.

```APIDOC
## Method: TmemBufferPool.allocate

### Description
Allocates a sub-region from the pool and returns a pointer. This method allocates a contiguous region of TMEM columns from the pool and returns a pointer to the start of that region.

### Parameters
- **size** (int | cutlass.cute.typing.Layout) - Required - The allocation size, which can be an explicit number of columns or a TMEM layout.
- **dtype** (Type) - Required - The data type for the returned pointer and for computing layout size.

### Returns
- **cute.Pointer** - A pointer to the allocated region with the specified dtype.

### Raises
- **AssertionError** - If there are not enough columns remaining in the pool.
```

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

### Get runtime fields of an Atom

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

Retrieves current runtime state values from an Atom instance.

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