NVIDIA CUTLASS

The CUTLASS Operator API provides a high-level Python interface for managing and integrating CUTLASS kernels.

Key Capabilities:

• Kernel Discovery: Use get_operators(args) to find kernels matching specific dtypes, layouts, and hardware constraints.
• Consistent Interface: Pass PyTorch or DLPack tensors directly into argument objects and call .run(args) without writing kernel-specific glue code.
• Custom Epilogue Fusions: Pass a plain Python function to fuse activations or elementwise operations into supported kernels via the CuTe DSL's Epilogue Fusion Configuration (EFC) framework.
• Bring-Your-Own-Kernel: Register custom CuTe DSL kernels to use them through the same interface as officially maintained kernels.
• Automatic Updates: Upgrading nvidia-cutlass-operators automatically provides access to new kernels, fixes, and optimizations in the registry.

CUTLASS 4.x provides Python-based Domain-Specific Languages (DSLs) to interface with the CUTLASS C++ template library. This allows for faster prototyping and easier iteration of high-performance linear algebra kernels on NVIDIA GPUs by using Python syntax instead of complex C++ metaprogramming.

The first DSL released is CuTe DSL, a low-level programming model consistent with CuTe C++ abstractions. It enables control over hardware threads, data hierarchy, and core concepts like layouts and tensors.

CuTe is a core library introduced in CUTLASS 3.0 for describing and manipulating tensors of threads and data. It uses C++ CUDA template abstractions to define hierarchically multidimensional layouts.

Key abstractions include:

• Layout: Compactly packages type, shape, memory space, and layout of data, performing complex indexing automatically.
• Tensor: Composes Layout objects with data arrays to represent tensors.

CuTe allows for functional composition of layouts, supporting operations like tiling and partitioning, which simplifies the design of dense linear algebra operations.

CUTLASS is a collection of CUDA C++ template abstractions designed for implementing high-performance matrix-matrix multiplication (GEMM) and related computations. It provides modular software components for hierarchical decomposition and data movement, allowing developers to specialize primitives via custom tiling sizes, data types, and algorithmic policies.

Key features include:

• Support for mixed-precision computations (FP64, FP32, TF32, FP16, BF16, 8b floating point, block scaled data types like NVFP4/MXFP4, and integer/binary types).
• Support across NVIDIA architectures: Volta, Turing, Ampere, Ada, Hopper, and Blackwell.
• CUTLASS DSLs: Python-native interfaces for writing high-performance CUDA kernels using CuTe concepts, offering faster compile times and easier integration with DL frameworks.

Blackwell (SM100) architecture introduces the tcgen05 family of PTX instructions, which are 5th-generation Tensor Core matrix multiply-accumulate (MMA) operations. These instructions compute D = A * B + C and offer 2x–4x the throughput of Hopper's WGMMA instructions depending on the data type.

Key architectural features include:

• Tensor Memory (TMEM): A dedicated on-chip memory for the accumulator (and optionally operand A). tcgen05 MMA reads/writes the accumulator directly in TMEM, reducing register file pressure.
• Single-thread launch: A single thread issues the MMA instruction.
• CTA-pair cooperation: Two adjacent CTAs can jointly execute a single MMA to double the tile size without additional synchronization logic.

The CUTLASS repository is organized into several key components:

• CUTLASS Template Library: Header-only CUDA C++ templates for linear algebra subroutines and solvers.
• CuTe Template Library: Core layout types and associated algebra (header-only).
• CUTLASS Instance Library: Procedurally generated instantiations of templates covering various configurations (data types, tile sizes, etc.).
• CUTLASS Profiler: A command-line application used to evaluate the functionality and performance of the Instance Library.
• CUTLASS Utilities: A companion library supporting tests, examples, and client applications.
• Examples: SDK examples demonstrating how to apply CUTLASS templates to basic computations.

In CuTe, specific ranks define logical structures:

• Vector: Any Layout with rank == 1. This includes hierarchical shapes that can be flattened into a single mode (e.g., ((4,2)):((2,1)) is a vector of size 8).
• Matrix: Any Layout with rank == 2.
  • Column-major: The first mode has a stride of 1.
  • Row-major: The second mode has a stride of 1.

Matrices can also have multi-mode dimensions (e.g., a column might be composed of multiple sub-modes), but they can still be indexed using 2-D coordinates.

The cute.arch module provides lightweight wrappers for NVVM Operation builders, allowing you to implement CUDA built-in device functions within the CuTe DSL. These wrappers integrate with CuTe DSL types and support source location tracking via the @dsl_user_op decorator.

Key capabilities include:

• Core CUDA built-ins: Access functions like thread_idx, warp_idx, block_dim, grid_dim, and cluster_dim.
• Memory Barrier Management: Manage barriers using functions such as mbarrier_init, mbarrier_arrive, and mbarrier_wait.
• Memory Allocation: Use SmemAllocator for low-level shared memory (SMEM) management and TmemAllocator for low-level tensor memory (TMEM) management.

A Layout is composed of a Shape and a Stride (both of which are IntTuple concepts). It maps coordinates within the Shape to an index via the Stride.

A Tensor is created by composing a Layout with data (such as a pointer or an array). The Layout generates the index used to subscript the data to retrieve elements.