CUTLASS Library: CUDA Templates for Linear Algebra

• CUTLASS is an open-source library for constructing highly optimized CUDA kernels for linear algebra operations on NVIDIA GPUs.
• It uses C++ template abstractions to enable flexible kernel composition, supporting fused operations and integration with deep learning compilers.
• The library achieves state-of-the-art performance through auto-tuning, efficient memory management, and GPU-specific optimizations.

The CUDA Templates for Linear Algebra Subroutines and Solvers (CUTLASS) library is an open-source collection and framework for constructing highly optimized CUDA C++ kernels for linear algebra operations—especially general matrix-matrix multiplication (GEMM) and its derivatives—on NVIDIA GPUs. CUTLASS is deeply architected around C++ templates and parametric composition, enabling algorithmic flexibility, high performance across GPU architectures (e.g., Ampere SM80, Hopper SM90), and modular extension to fused operations. As both a low-level back-end for commercial and research deep learning compilers and a direct target for specialist kernel programmers, CUTLASS encodes expert knowledge of GPU microarchitecture, memory hierarchies, and kernel fusion strategies, making it a central asset for software-hardware codesign in high-performance deep learning workloads (Bikshandi et al., 2023, Heidari et al., 29 Apr 2026).

1. Architecture and Abstraction Model

CUTLASS is organized around a hierarchy of C++ template abstractions that define key properties of data types, tiling strategies, threadblock-level workloads, warp-level microkernels, and epilogue operations. These templates enable instantiation of device kernels that match the wide combinatorial space of problem shapes, memory layouts, data types (fp32, fp16, bf16, tf32), and hardware-specific execution strategies.

The central abstractions in CUTLASS include:

• Layouts and Tensors: Layouts (e.g., cutlass::layout::RowMajor, ColMajor, and Hopper-specific swizzled/composed layouts) define the mapping of logical $\mathrm{N}$-dimensional tensor coordinates to linear memory. Tensors abstract typed memory pointers and their associated layouts (Bikshandi et al., 2023).
• GEMM Templates: At the core is cutlass::gemm::device::Gemm, which is parameterized by data types $A$, $B$, $C$, memory layouts, hardware operation class (OpClassTensorOp, OpClassSimt), streaming multiprocessor architecture tag (e.g., Sm80, Sm90), and tiling strategies (threadblock, warp, instruction tile shapes).
• Epilogue Functors: These define post-processing operations (e.g., scaling, activation, bias, softmax) fused into the output stage of the kernel—often key in transformer and DL workloads (Heidari et al., 29 Apr 2026).

This compositional model enables both library builders and code generation systems to programmatically instantiate kernels with precise architectural and algorithmic specializations.

2. Kernel Generation and Tuning Workflow

CUTLASS kernel synthesis proceeds through a multi-stage pipeline, often orchestrated by higher-level workflows such as FACT (Framework for Agentic CUTLASS Transpilation) (Heidari et al., 29 Apr 2026):

1. Graph Tracing and Pattern Matching: Leveraging PyTorch torch.fx or torch.jit, the upstream system emits a computational graph, which is then analyzed for subgraphs matching known optimization patterns (e.g., GEMM, fused MLP, FMHA). Each match records shapes, dtypes, and target architecture.
2. Template Parameterization: Extracted subgraph properties are used to select or generate explicit C++ template arguments: GEMM rule $\to$ kernel variant; $M,N,K$ (problem shapes) $\to$ tile sizes; dtypes $\to$ element types and scaling; architecture $\to$ specific tag (e.g., Sm80, Sm90). This forms an instantiation spec lifted into C++ templates.
3. Vetted Example Index: A registry, indexed by pattern/dtype/arch, holds curated example CUTLASS instantiations (recording sample paths, template parameters, scheduling modes, and epilogue options), supporting nearest-neighbor retrieval to ground new kernel synthesis (see Table 1).

Key	Example Fields	Role
rule, dtype, arch	sample path, tile sizes, schedule, epilogue	Indexing/lifting instantiations

1. Kernel Synthesis and Auto-Tuning: New kernels are emitted as C++ extensions (e.g., for PyTorch via torch.utils.cpp_extension.load), compiled, functionally verified against reference ops (e.g., torch.matmul), and empirically tuned by grid search over tile sizes, pipeline stages, scheduling options. Device constraints (shared memory, registers) prune the search space. Performance is recorded using metrics such as average FLOPS and latency.
2. End-to-End Integration: Final kernel modules register as Python-callable ops, replacing traced graph submodules for deployment (Heidari et al., 29 Apr 2026).

3. Memory and Execution Model on Recent Architectures

CUTLASS is highly tuned to leverage GPU architectural features. On NVIDIA Hopper (SM90), key hardware primitives exposed through CUTLASS include:

• Tensor Memory Accelerator (TMA): Enables asynchronous, 128-byte-aligned DMA from HBM to SMEM, supporting efficient tiling and overlap of copy/compute.
• Warpgroup MMA (WGMMA): 128-thread, high-throughput tensor instructions effectively doubling FP16 throughput relative to SM80 cores (Bikshandi et al., 2023).
• Tile and Copy Abstractions: Layouts, Tensors, and CuTe helpers enable mapping of logical operations to swizzled shared memory tiles, minimizing bank conflicts.

In the case of fused attention kernels (e.g., FlashAttention-2), CUTLASS expresses the entire $QK^T \to \text{softmax} \to (P) \to PV$ chain as a single fully fused kernel:

• Tiles for $A$0, $A$1, $A$2 are defined in global and shared memory via layout composition and tiling.
• Online softmax and memory-efficient attention mechanisms keep all intermediate state in on-chip registers (rmem), never materializing large $A$3 matrices to DRAM.
• Kernel execution interleaves asynchronous TMA copy, WGMMA GEMM execution, and fused epilogue (softmax, dropout, bias) in threadblocks, optimizing arithmetic intensity and register/shmem utilization.
• Hardware occupancy (register pressure, shared memory per CTA) is tuned by empirical sweep of tile sizes per head dimension. Empirical results show configurations such as $A$4 or $A$5 optimize for multi-head attention workloads (Bikshandi et al., 2023).

4. Support for Fused and Compositional Kernels

CUTLASS modularity enables construction of complex fused kernels:

• Fused Multi-Head Attention (FMHA): CUTLASS instantiates custom epilogue functors to express online softmax, bias addition, and dropout mask application.
• Fused MLP blocks: Instantiated as sequences of back-to-back GEMMs fused with elementwise activation epiogues (e.g., GELU, SiLU).
• Template-based kernel composition: Complex fused operations are constructed by composing and specializing CUTLASS primitives, maintaining high arithmetic throughput and minimizing DRAM traffic (Heidari et al., 29 Apr 2026).

This extensibility is a key enabler for agent-driven or LLM-based kernel synthesis workflows, which can ground new variants in a library of vetted kernel patterns.

5. Benchmarking and Empirical Results

CUTLASS-based kernels have been benchmarked extensively against vendor libraries (e.g., cuBLAS) and stock deep learning compilers (e.g., PyTorch eager, Inductor, TensorRT). Key results include:

• On A100 (Ampere SM80) GPUs, auto-tuned CUTLASS GEMM kernels achieve 1.06x–1.18x speedup over cuBLAS (Level 1 GEMM, square/batched/large-K).
• On H100 (Hopper SM90), performance varies 0.84x–1.80x relative to cuBLAS, reflecting sensitivity to microarchitecture and tile selection.
• For Level 3 transformer blocks, the FACT workflow achieves 2.03x speedup on models like MiniGPT against PyTorch eager (compared to Inductor: 1.89x, TensorRT: 1.85x), and 1.41x on Llama 3 8B (Heidari et al., 29 Apr 2026).
• For FlashAttention-2 forward pass, a CUTLASS/Hopper implementation achieved 20–50% higher TFLOPs/s over Ampere-optimized baselines and ~2.5–3× the throughput of stock CUTLASS FMHA example kernels, primarily due to the utilization of TMA, WGMMA instructions, and kernel fusion strategies (Bikshandi et al., 2023).

6. Best Practices and Limitations

Achieving peak performance with CUTLASS requires:

• Careful parameterization: aligning problem sizes ($A$6) to multiples of tile dimensions and hardware warp sizes.
• Device-aware auto-tuning: ensuring no configuration exceeds hardware shared memory (sharedMemPerBlock) or register (regsPerBlock) limits; pruning search space with a simple roofline cost model.
• Warmup runs and runtime benchmarking for valid performance evaluation.
• Caching compiled kernel objects keyed by template hashes to avoid redundant builds.
• Verifying numerical equivalence, particularly when fusing epilogues or using mixed precision; tolerance thresholds should account for datatype-induced discrepancies.
• Leveraging CUTLASS status/debug facilities to diagnose template instantiation and runtime mismatches (Heidari et al., 29 Apr 2026).

A plausible implication is that while CUTLASS provides a high-level compositional approach, realizing its full potential in new architectures or workloads typically requires a meta-workflow (such as FACT) that automates template discovery, instantiation, and empirical validation.

7. Significance in Deep Learning Systems

CUTLASS is a foundational kernel library underlying both hand-tuned research kernels and agentic/automated compilation workflows. It abstracts critical microarchitectural details, provides a robust substrate for high-throughput matrix and attention primitives, and integrates with compiler-based and LLM-driven code synthesis pipelines (Heidari et al., 29 Apr 2026, Bikshandi et al., 2023). By enabling direct expression of fused patterns, asynchronous memory transfers, and specialized epilogues, CUTLASS supports rapid adaptation to new GPU architectures and evolving deep learning model structures.