| name | cutlass |
| description | NVIDIA CUTLASS CUDA Template Library - comprehensive reference for high-performance matrix multiplication (GEMM), convolution, tensor operations, and CuTe DSL across all GPU architectures (Volta through Blackwell) |
| version | 3.8 |
CUTLASS - CUDA Template Library for Linear Algebra
CUTLASS (CUDA Templates for Linear Algebra Subroutines) is NVIDIA's open-source C++ template library for implementing high-performance matrix-matrix multiplication (GEMM) and related computations at all levels and scales within CUDA. It provides near-optimal utilization of peak theoretical throughput on NVIDIA GPUs.
Overview
CUTLASS provides:
- High-performance GEMM operations with near-optimal utilization of peak theoretical throughput
- Support for multiple NVIDIA architectures: Volta (SM70), Turing (SM75), Ampere (SM80/SM89), Ada (SM89), Hopper (SM90), and Blackwell (SM100/SM101/SM103/SM120)
- Extensive data type support: FP64, FP32, TF32, FP16, BF16, FP8 (e5m2, e4m3), INT8, INT4, INT2, binary 1b, and block-scaled types (NVFP4, MXFP4/6/8)
- CUTLASS DSLs: Python native interfaces for writing high-performance CUDA kernels (CuTe DSL)
- CuTe library: A modern C++ tensor abstraction for composable GPU micro-kernels
When to Use This Skill
Use this skill when:
- Writing or optimizing GEMM kernels using CUTLASS
- Implementing convolution operations with CUTLASS
- Using CuTe for tensor operations and layout algebra
- Working with Tensor Core operations on NVIDIA GPUs
- Implementing mixed-precision training or inference kernels
- Building custom epilogue fusion kernels
- Using TMA (Tensor Memory Accelerator) on Hopper/Blackwell
- Implementing fused attention (FMHA) kernels
- Working with sparse or block-scaled GEMM operations
- Profiling and benchmarking CUTLASS kernels
Quick Reference
CUTLASS 3.x GEMM (Recommended)
#include "cutlass/cutlass.h"
#include "cutlass/gemm/device/gemm_universal_adapter.h"
#include "cutlass/gemm/kernel/gemm_universal.h"
#include "cutlass/epilogue/collective/default_epilogue.hpp"
#include "cutlass/gemm/collective/collective_builder.hpp"
using ElementA = cutlass::half_t;
using ElementB = cutlass::half_t;
using ElementC = float;
using ElementD = cutlass::half_t;
using LayoutA = cutlass::layout::RowMajor;
using LayoutB = cutlass::layout::ColumnMajor;
using LayoutC = cutlass::layout::RowMajor;
using LayoutD = cutlass::layout::RowMajor;
using CollectiveOp = typename cutlass::gemm::collective::CollectiveBuilder<
cutlass::arch::Sm90, cutlass::arch::OpClassTensorOp,
ElementA, LayoutA, 8,
ElementB, LayoutB, 8,
ElementC,
cutlass::gemm::GemmShape<128, 128, 64>,
cutlass::gemm::collective::StageCountAutoCarveout<0>,
cutlass::gemm::collective::KernelScheduleAuto
>::CollectiveOp;
using EpilogueOp = cutlass::epilogue::collective::DefaultEpilogue<
LayoutD, LayoutC,
cutlass::epilogue::collective::EpilogueScheduleAuto
>;
using Kernel = cutlass::gemm::kernel::GemmUniversal<
cutlass::gemm::collective::CollectiveOp,
cutlass::epilogue::collective::EpilogueOp
>;
using Gemm = cutlass::gemm::device::GemmUniversalAdapter<Kernel>;
Gemm gemm_op;
auto args = Gemm::Arguments{
{M, N, K},
{ptr_A, stride_A},
{ptr_B, stride_B},
{ptr_C, stride_C},
{ptr_D, stride_D},
{alpha, beta}
};
gemm_op(args);
CUTLASS 2.x GEMM
#include "cutlass/gemm/device/gemm.h"
using Gemm = cutlass::gemm::device::Gemm<
cutlass::half_t, cutlass::layout::RowMajor,
cutlass::half_t, cutlass::layout::ColumnMajor,
float, cutlass::layout::RowMajor,
float,
cutlass::arch::OpClassTensorOp,
cutlass::arch::Sm80,
cutlass::gemm::GemmShape<128, 128, 32>,
cutlass::gemm::GemmShape<64, 64, 32>,
cutlass::gemm::GemmShape<16, 8, 16>,
cutlass::epilogue::thread::LinearCombination<
float, 4, float, float>
>;
Gemm gemm_op;
gemm_op({M, N, K}, alpha, ptr_A, lda, ptr_B, ldb, beta, ptr_C, ldc, ptr_D, ldd);
Architecture Support
| Architecture | SM Version | Key Features |
|---|
| Volta | SM70 | First Tensor Cores (FP16), WMMA API |
| Turing | SM75 | Tensor Cores (INT8/INT4/INT1), MMA instructions |
| Ampere | SM80/SM89 | TF32, BF16, async copy, cp.async |
| Hopper | SM90 | TMA, GMMA, WGMMA, thread block clusters, warp specialization |
| Blackwell | SM100+ | UMMA, block-scaled types (NVFP4/MXFP), distributed GEMM, green contexts |
Key Concepts
- GEMM Hierarchy: Device > Kernel > Collective > Tiled MMA/Copy > Atom
- Collective: Mainloop computation shared across threadblock/warp
- Epilogue: Post-GEMM operations (activation, scaling, type conversion)
- CuTe Layout: Formalized layout algebra for thread-to-data mapping
- TensorRef/TensorView: Lightweight tensor abstractions with layout
- Pipeline: Multi-stage data movement with producer-consumer pattern
- Split-K: Parallel reduction across K dimension
- Warp Specialization: Separate producer and consumer warp groups (SM90+)
Reference Chapters
Detailed reference documentation is organized in the references/ directory:
Fundamentals
- 01-overview-and-architecture.md - Project overview, design philosophy, and architecture
- 02-quickstart.md - Prerequisites, build instructions, first kernel
- 03-code-organization.md - Directory structure and code layout
- 04-data-types.md - Numeric types, conversions, and type traits
- 05-layout-system.md - Matrix/tensor layout definitions (RowMajor, ColumnMajor, etc.)
- 06-tensor-abstractions.md - TensorRef, TensorView, coordinate systems
- 38-terminology.md - Comprehensive glossary of CUTLASS terminology
GEMM API
Post-Processing
Data Movement
Operations
CuTe Library
Advanced Topics
Architecture-Specific
Tools and Utilities
