> ## Documentation Index
> Fetch the complete documentation index at: https://mintlify.com/NVIDIA/cutlass/llms.txt
> Use this file to discover all available pages before exploring further.

# C++ API Overview

> Overview of the CUTLASS C++ template library and core APIs

## Introduction

CUTLASS is a collection of CUDA C++ template abstractions for implementing high-performance matrix-matrix multiplication (GEMM), convolution, and related computations at all levels and scales within CUDA. It incorporates strategies for hierarchical decomposition and data movement similar to those used to implement cuBLAS and cuDNN.

## Core Components

The CUTLASS C++ API is organized into several key namespaces and components:

### Namespace Organization

<CodeGroup>
  ```cpp Main Namespaces theme={null}
  cutlass::                   // Top-level namespace
  cutlass::gemm::             // GEMM operations
  cutlass::conv::             // Convolution operations
  cutlass::epilogue::         // Epilogue operations
  cutlass::arch::             // Architecture-specific components
  cute::                      // CuTe library (modern tensor abstraction)
  ```
</CodeGroup>

### Device-Level APIs

CUTLASS provides device-level operator templates that can be instantiated and launched from host code:

<ParamField path="cutlass::gemm::device::Gemm" type="template class">
  Primary GEMM device operator for matrix multiplication
</ParamField>

<ParamField path="cutlass::gemm::device::GemmUniversal" type="template class">
  Universal GEMM supporting batched, grouped, and split-K modes
</ParamField>

<ParamField path="cutlass::conv::device::ImplicitGemmConvolution" type="template class">
  Implicit GEMM convolution device operator
</ParamField>

## Template Design Pattern

CUTLASS uses extensive compile-time template metaprogramming to optimize kernels:

<CodeGroup>
  ```cpp Type Definitions theme={null}
  // Element types
  using ElementA = float;
  using ElementB = float;
  using ElementC = float;
  using ElementAccumulator = float;

  // Layout types
  using LayoutA = cutlass::layout::ColumnMajor;
  using LayoutB = cutlass::layout::ColumnMajor;
  using LayoutC = cutlass::layout::ColumnMajor;

  // Architecture tag
  using ArchTag = cutlass::arch::Sm80;

  // Operator class
  using OperatorClass = cutlass::arch::OpClassTensorOp;
  ```

  ```cpp Complete Instantiation theme={null}
  using GemmKernel = cutlass::gemm::device::Gemm<
    float,                              // ElementA
    cutlass::layout::ColumnMajor,       // LayoutA
    float,                              // ElementB
    cutlass::layout::ColumnMajor,       // LayoutB
    float,                              // ElementC
    cutlass::layout::ColumnMajor        // LayoutC
  >;
  ```
</CodeGroup>

## Hierarchical Decomposition

CUTLASS organizes computation at multiple levels:

<CardGroup cols={3}>
  <Card title="Device Level" icon="microchip">
    Kernel launch and grid management
  </Card>

  <Card title="Threadblock Level" icon="grid">
    Cooperative thread array (CTA) tiling
  </Card>

  <Card title="Warp Level" icon="layer-group">
    Warp-level matrix operations
  </Card>

  <Card title="Thread Level" icon="code">
    Per-thread computation and data movement
  </Card>

  <Card title="Instruction Level" icon="bolt">
    Hardware-accelerated instructions (Tensor Cores)
  </Card>
</CardGroup>

## Tile Shapes

Performance is controlled by tile sizes at each level:

<ParamField path="ThreadblockShape" type="GemmShape">
  Threadblock-level tile size (e.g., `GemmShape<128, 128, 32>`)
</ParamField>

<ParamField path="WarpShape" type="GemmShape">
  Warp-level tile size (e.g., `GemmShape<64, 64, 32>`)
</ParamField>

<ParamField path="InstructionShape" type="GemmShape">
  Instruction-level tile size (e.g., `GemmShape<16, 8, 16>` for Tensor Cores)
</ParamField>

## Common Data Types

CUTLASS defines several core data structures:

### GemmCoord

```cpp theme={null}
cutlass::gemm::GemmCoord problem_size(M, N, K);
```

Represents the dimensions of a GEMM problem (M×N = M×K \* K×N).

### TensorRef

```cpp theme={null}
cutlass::TensorRef<ElementType, LayoutType> tensor_ref(ptr, leading_dimension);
```

A lightweight reference to a tensor in memory with layout information.

### Status

```cpp theme={null}
cutlass::Status status = kernel(args);
if (status != cutlass::Status::kSuccess) {
  // Handle error
}
```

Return codes for CUTLASS operations.

## API Usage Pattern

The typical workflow for using CUTLASS device operators:

<Steps>
  <Step title="Define the kernel type">
    Instantiate a CUTLASS template with desired types and parameters
  </Step>

  <Step title="Create kernel instance">
    Construct an instance of the kernel operator
  </Step>

  <Step title="Prepare arguments">
    Construct an Arguments structure with problem size, tensor references, and parameters
  </Step>

  <Step title="Check feasibility">
    Call `can_implement()` to verify the kernel can execute the problem
  </Step>

  <Step title="Get workspace size">
    Call `get_workspace_size()` if the kernel requires temporary storage
  </Step>

  <Step title="Initialize kernel">
    Call `initialize()` with arguments and workspace
  </Step>

  <Step title="Execute kernel">
    Call `run()` or the function call operator to launch the kernel
  </Step>
</Steps>

<Note>
  The `initialize()` and `run()` steps can be combined using the function call operator: `status = kernel(args, workspace, stream);`
</Note>

## File Organization

Key header files in the CUTLASS C++ API:

```
include/cutlass/
├── cutlass.h                    # Core definitions
├── gemm/
│   ├── gemm.h                  # GEMM type definitions
│   ├── device/
│   │   ├── gemm.h              # Device-level GEMM
│   │   └── gemm_universal.h    # Universal GEMM
│   └── kernel/                 # Kernel implementations
├── conv/
│   ├── convolution.h           # Convolution types
│   └── device/                 # Device-level convolution
├── epilogue/
│   └── thread/                 # Epilogue operations
└── layout/
    └── matrix.h                # Layout types

include/cute/
├── tensor.hpp                  # CuTe tensor abstraction
├── layout.hpp                  # CuTe layouts
└── algorithm/                  # CuTe algorithms
```

## Architecture Support

CUTLASS supports multiple NVIDIA GPU architectures:

<CodeGroup>
  ```cpp Volta (SM70) theme={null}
  cutlass::arch::Sm70
  ```

  ```cpp Turing (SM75) theme={null}
  cutlass::arch::Sm75
  ```

  ```cpp Ampere (SM80) theme={null}
  cutlass::arch::Sm80
  ```

  ```cpp Hopper (SM90) theme={null}
  cutlass::arch::Sm90
  ```

  ```cpp Blackwell (SM100) theme={null}
  cutlass::arch::Sm100
  ```
</CodeGroup>

## Operator Classes

Different compute capabilities:

<ParamField path="cutlass::arch::OpClassSimt" type="tag">
  CUDA cores (SIMT) - available on all architectures
</ParamField>

<ParamField path="cutlass::arch::OpClassTensorOp" type="tag">
  Tensor Cores - available on Volta and later
</ParamField>

<ParamField path="cutlass::arch::OpClassWmmaTensorOp" type="tag">
  WMMA Tensor Cores - Volta-specific
</ParamField>

## Next Steps

<CardGroup cols={2}>
  <Card title="GEMM API" icon="calculator" href="/cpp/gemm-api">
    Learn about the GEMM API and usage patterns
  </Card>

  <Card title="CuTe Library" icon="cube" href="/cpp/cute-library">
    Explore the modern CuTe tensor abstraction
  </Card>

  <Card title="Epilogue Operations" icon="wand-magic-sparkles" href="/cpp/epilogue">
    Add custom operations to your kernels
  </Card>

  <Card title="Convolution API" icon="grip" href="/cpp/convolution">
    Perform efficient convolution operations
  </Card>
</CardGroup>