Engineering

CUTLASS vs Triton Kernel Authoring: CuTe DSL Guide (2026)

CUTLASS vs Triton Kernel AuthoringCuTe DSLCUTLASS 4.4GB300Blackwell Kernel DevelopmentGPU Kernel DevelopmentH100B200GPU Cloud
CUTLASS vs Triton Kernel Authoring: CuTe DSL Guide (2026)

CUTLASS vs Triton kernel authoring is the question NVIDIA's own changelog never quite answers, and CUTLASS's CuTe DSL beta status makes this the right moment to settle it directly. CuTe DSL has been in public beta since CUTLASS 4.0 shipped its Python front end, and NVIDIA's own FAQ page commits to graduating it out of beta by the end of summer 2026. That window has effectively closed as of this post's publish date, which makes now the right moment to ask what the changelog does not spell out: for kernel authoring on H100 and Blackwell, is CuTe DSL worth learning over Triton, and where does each one actually win. This post works through a real GEMM in CuTe DSL, compares it against cuBLAS, and ends with a decision framework instead of a changelog recap.

If you have not written a custom GPU kernel before, start with our OpenAI Triton kernel development guide for the fundamentals of tile-level programming in Python; this post assumes that background and focuses on where CUTLASS's CuTe DSL diverges from it.

CUTLASS vs Triton Kernel Authoring: What Changed with CuTe DSL's Path to GA

CuTe DSL is CUTLASS's Python-embedded kernel language, distinct from Triton and distinct from CUDA 13's own cuTile Python DSL, and its beta status is about to matter for anyone deciding which tool to standardize on this year.

From CUTLASS 4.0's Python Front End to a Public Beta Nearing GA

CUTLASS has shipped C++ template kernels for GEMM and convolution since its earliest releases, but writing against those templates required real C++ template metaprogramming skill: specializing tile shapes, epilogues, and warp layouts through nested template parameters. CUTLASS 4.0 added CuTe DSL as a Python front end onto the same underlying tile abstractions (CuTe, the "CUDA Templates" layout algebra) that the C++ library uses internally. Instead of instantiating C++ templates, you write a Python function decorated for JIT compilation, and it lowers through the same MLIR-based pipeline as the C++ path.

NVIDIA's CuTe DSL FAQ states plainly that the DSL is currently in public beta and that graduation out of beta is targeted for the end of summer 2026. That is a specific, dated commitment from NVIDIA itself, not an inferred roadmap item, which is part of why the timing of this post matters: teams evaluating CuTe DSL for anything beyond experimentation have been waiting on exactly this milestone.

CUTLASS 4.4: CTK 13.1 and GB300/Blackwell Ultra Support

CUTLASS 4.4.0 landed CUDA Toolkit 13.1 support and, with it, made GB300 (Blackwell Ultra, compute capability SM103) a working CuTe DSL target under CTK 13.1, including a new SM103 batched 3xFP4 blockscaled GEMM example shipped in the same release. That is a meaningful jump: SM103 is a newer, narrower target than the SM100 (B200) support CUTLASS already had, and blockscaled FP4 GEMM is exactly the kind of low-precision, high-throughput kernel that makes hand-authoring worth the effort in the first place.

More broadly, CuTe DSL's supported architecture range runs from Ampere (SM80) through Blackwell. That range is worth noting precisely because it excludes nothing newer than what CUTLASS's C++ path already covers: CuTe DSL is not a reduced-capability preview, it is the same architecture coverage in a different front end.

CUTLASS vs Triton Kernel Authoring: Where Each One Wins

Neither tool replaces the other across the board, and the honest answer to "CUTLASS or Triton" is "it depends which operation you are authoring."

Where Triton Still Wins: Pointwise Ops, Reductions, Anything Memory-Bound

Triton's tile-level abstraction (tl.load, tl.store, block pointers, autotuning) is the faster path for elementwise and reduction kernels: fused softmax, layer norm, RMS norm, activation fusions, and anything where the bottleneck is HBM bandwidth rather than tensor-core throughput. Triton's autotuner sweeps block size and warp count automatically, and the ecosystem already assumes Triton: torch.compile's inductor backend emits Triton by default, and vLLM's PagedAttention, RoPE, and RMS Norm kernels are Triton under the hood. Our Triton kernel development guide walks through a fused softmax kernel end to end, including the persistent-kernel pattern that keeps SMs busy on memory-bound work.

If your custom operation is not a matrix multiply and does not need warp-specialized tensor-core scheduling, Triton is very likely still the right tool. It is also the tool with the shallower learning curve: no CuTe layout algebra to learn, no separate tile-and-fragment mental model.

Where CuTe DSL Wins: GEMMs, Attention, and the Warp-Specialized Blackwell Path

CuTe DSL's advantage shows up specifically on operations that are compute-bound on tensor cores and benefit from warp-specialized scheduling: GEMM, grouped GEMM, and fused attention. NVIDIA's own benchmark of dense GEMM in CuTe DSL against the C++ CUTLASS implementation found that CuTe DSL matches C++ CUTLASS throughput on Blackwell across most problem sizes, with one acknowledged gap at small K-sizes (K=512) where synchronization overhead currently costs it performance; on Ampere/A100, CuTe DSL dense GEMM is still slightly slower than the C++ path, an open gap NVIDIA has not closed yet. That is a strong claim to make about a Python DSL: matching hand-written C++ template performance on a GEMM is exactly the kind of parity Triton has never claimed for itself on the same operation class.

The clearest production proof point is FlashAttention-4, the reference attention kernel for Blackwell. FA4 was authored by Tri Dao and collaborators from Princeton, Together AI, Meta, NVIDIA, and Colfax, and it is written entirely in CuTe DSL, a Python DSL rather than C++. That last distinction is worth being precise about, because it is easy to conflate with two other, similarly-named things: CUDA 13's own cuTile Python DSL is a separate NVIDIA project, not part of CUTLASS, and it targets a different abstraction layer; our CUDA 13 tile programming guide covers that side of the landscape. CuTe DSL and cuTile share a "Python, tile-level, NVIDIA-authored" description, and nothing else. For the attention kernel itself, see our FlashAttention-4 Blackwell inference guide for setup and serving benchmarks.

On raw numbers, Lambda's own benchmark of FA4 on B200 reports 1,613 TFLOPs/s in BF16, 71% of peak hardware utilization, 1.3x faster than cuDNN 9.13 and 2.7x faster than a Triton implementation of the same kernel. That last comparison is the sharpest data point in this whole post: on the specific operation of warp-specialized fused attention on Blackwell, a CuTe DSL kernel outran a Triton kernel by nearly 3x. That gap is why FA4 was not written in Triton, and it is the strongest concrete argument for learning CuTe DSL if attention or GEMM authoring is actually your bottleneck.

Raw CUDA C++: Still the Floor Both DSLs Compile Down To

Neither Triton nor CuTe DSL eliminates CUDA C++, they both sit on top of it. Triton's compiler emits PTX/CUBIN through its own MLIR-based pipeline; CuTe DSL emits through the same underlying tile abstractions the C++ CUTLASS templates use. When a kernel needs a hardware feature neither DSL has exposed yet, or when debugging needs to drop to PTX-level inspection, C++ CUTLASS or hand-written CUDA remains the actual floor. For most teams that floor is now rarely visited: CuTe DSL's documented limitations list is short enough that the gap versus C++ CUTLASS is a known, bounded list rather than an open-ended one (see the debugging section below).

A GEMM Kernel in CuTe DSL, Walked Through

The fastest way to evaluate CuTe DSL against Triton for your own use case is to write the same kernel in both and compare. Here is a batched GEMM in CuTe DSL, walked through piece by piece.

Environment and Version Pinning

Given the portability caveat above, pin every layer explicitly rather than floating on pip install cutlass:

bash
# CTK 13.1 is required for GB300/SM103 targets in CUTLASS 4.4.x;
# CTK 12.8+ is sufficient for Ampere through B200 (SM80-SM100).
pip install nvidia-cutlass-dsl==4.4.0
pip install torch --index-url https://download.pytorch.org/whl/cu128

python -c "import cutlass; print(cutlass.__version__)"
nvidia-smi --query-gpu=name,compute_cap --format=csv,noheader

Record the exact cutlass and torch versions in your Dockerfile or lockfile. Because CuTe DSL makes no cross-release portability guarantee during beta, a kernel that compiles today against 4.4.0 is not guaranteed to compile unchanged against 4.5.0. Treat every version bump as a re-validation event, not a drop-in upgrade, the same discipline that matters for pinning triton against the exact torch build that bundles it.

The Kernel: Tiles, TMA Loads, and the Fragment-Free API

CuTe DSL's GEMM kernel expresses the problem as a set of tiles moved through the memory hierarchy via TMA (Tensor Memory Accelerator) loads, rather than as explicit thread-indexed loops. This is the biggest conceptual shift coming from Triton: Triton's tl.make_block_ptr still has you reasoning about a single 2D tile per program instance, while CuTe DSL's layout algebra lets you describe multi-dimensional tile hierarchies (thread block tile, warp tile, instruction tile) in one composed layout.

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

@cute.kernel
def gemm_kernel(
    gA: cute.Tensor, gB: cute.Tensor, gC: cute.Tensor,
    tiled_mma: cute.TiledMma,
    tma_atom_a: cute.CopyAtom, tma_atom_b: cute.CopyAtom,
):
    tidx, _, _ = cute.arch.thread_idx()
    bidx, bidy, _ = cute.arch.block_idx()

    # Slice the global tensors down to this thread block's tile via TMA.
    # TMA loads move a full tile HBM -> shared memory in one async op,
    # with no manual address computation the way Triton's tl.load needs.
    smem_a = cute.make_fragment(tma_atom_a.shape, cutlass.Float16)
    smem_b = cute.make_fragment(tma_atom_b.shape, cutlass.Float16)

    cute.copy(tma_atom_a, gA[(None, None, bidx)], smem_a)
    cute.copy(tma_atom_b, gB[(None, None, bidy)], smem_b)
    cute.arch.cp_async_wait_group(0)
    cute.arch.barrier()

    # The MMA atom issues warp-group-level tensor core instructions
    # directly against the shared-memory tiles; there is no separate
    # "load into registers as fragments" step to author by hand.
    acc = cute.make_fragment(tiled_mma.shape_mnk[:2], cutlass.Float32)
    cute.gemm(tiled_mma, acc, smem_a, smem_b, acc)

    cute.copy(cute.make_copy_atom_c(), acc, gC[(None, None, bidx, bidy)])


@cute.jit
def gemm(A: cute.Tensor, B: cute.Tensor, C: cute.Tensor):
    tiled_mma = cute.make_tiled_mma(cute.arch.mma_op_f16_f32())
    tma_atom_a, tma_tensor_a = cute.make_tma_atom(A, tiled_mma)
    tma_atom_b, tma_tensor_b = cute.make_tma_atom(B, tiled_mma)
    gemm_kernel(tma_tensor_a, tma_tensor_b, C, tiled_mma, tma_atom_a, tma_atom_b).launch(
        grid=(A.shape[0] // 128, B.shape[1] // 128, 1),
        block=(128, 1, 1),
    )


def run_gemm(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
    c = torch.empty(a.shape[0], b.shape[1], device="cuda", dtype=torch.float32)
    gemm(from_dlpack(a), from_dlpack(b), from_dlpack(c))
    return c

Two things stand out against the Triton equivalent. First, there is no BLOCK_K loop written by hand: the TMA copy atom and the tiled MMA description carry the K-dimension iteration inside their layout, so the kernel body reads closer to "describe the tiles, issue the copy, issue the matmul" than "loop over K and accumulate." Second, there is no separate fragment-loading step the way raw CUDA C++ WMMA/wgmma code needs; cute.gemm issues the tensor-core instruction directly against the shared-memory tile description. This is what NVIDIA means when it says CuTe DSL targets C++ CUTLASS performance from a Python front end: the abstraction is higher-level than Triton's block pointers, but the generated instruction stream still goes through the same warp-specialized tensor-core path as hand-written CUTLASS.

What Breaks First: Debugging Without a C++-Grade Toolchain

The honest cost of CuTe DSL today is debugging. NVIDIA's own limitations documentation is direct about this: debugging tooling for CuTe DSL is "more limited in comparison to the C++ API," and there is no single-stepping. In practice, that means a layout mismatch (a TMA atom shape that does not match the tensor you fed it) surfaces as a runtime shape-assertion failure rather than something a debugger lets you step through instruction by instruction.

The same limitations page lists three other constraints worth knowing before you commit real kernel work to CuTe DSL: it has no convolution support today, layouts are restricted to 32-bit shapes and strides, and the DSL does not support dependent types or early returns inside control flow. None of these are showstoppers for GEMM or attention authoring, the two operation classes CuTe DSL is actually built for, but they rule out a direct port of anything convolution-based, and they mean control flow inside a kernel body has to be structured more conservatively than equivalent Python or Triton code.

Practical debugging workflow that holds up given these limits: validate correctness against a PyTorch/cuBLAS reference on small shapes first (torch.allclose with a loose tolerance for FP16/BF16 accumulation differences), then move to Nsight Compute for performance debugging once correctness is established. ncu's roofline and Memory Workload Analysis views work identically against CuTe DSL kernels as against C++ CUTLASS or Triton kernels, since they profile the compiled CUBIN rather than the source language. That profiling step needs the same root-level hardware counter access that bare-metal instances provide and that shared, virtualized cloud VMs frequently block.

A Benchmark Protocol: CuTe DSL vs cuBLAS on a Custom Fused Op

The Op and the Hardware

Here is a reproducible protocol for this comparison, built to run on Spheron bare-metal H100 SXM5 and B200 SXM6 instances, targeting a fused GEMM-plus-bias-epilogue op: a matmul immediately followed by a bias add and a GELU activation, folded into the GEMM's epilogue instead of run as separate kernels afterward. This is a realistic stand-in for the kind of op teams actually reach for custom kernel authoring on: it is common in MLP blocks, it is compute-bound enough to benefit from tensor-core scheduling, and cuBLAS alone cannot fuse the epilogue for you, since cuBLAS is a pure GEMM library with no activation-fusion hooks.

To run the same comparison yourself, start with the plain-GEMM baseline below: it benchmarks run_gemm, the CuTe DSL kernel defined above, against a plain torch.matmul call using CUDA events. Get this baseline running and correct first, then add the bias-add-plus-GELU epilogue to both sides, F.gelu(torch.matmul(a, b) + bias) for the cuBLAS side, and a GELU epilogue in the cute.copy step to the output tensor for the CuTe DSL side, before drawing any conclusion about the fused case:

python
import torch
from cutlass.cute.runtime import from_dlpack

a = torch.randn(4096, 4096, device="cuda", dtype=torch.float16)
b = torch.randn(4096, 4096, device="cuda", dtype=torch.float16)

def cublas_baseline(a, b):
    return torch.matmul(a, b)

start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
for fn, label in [(cublas_baseline, "cuBLAS (torch.matmul)"), (run_gemm, "CuTe DSL (run_gemm)")]:
    for _ in range(10):
        fn(a, b)  # warm up
    start.record()
    for _ in range(100):
        fn(a, b)
    end.record()
    torch.cuda.synchronize()
    print(f"{label}: {start.elapsed_time(end) / 100:.3f} ms")

What to Expect: No Sweep Has Been Run Yet, Here Is the Closest Published Anchor

This specific fused op has not been benchmarked across a full shape sweep by NVIDIA or by us; the protocol above is something to run yourself before trusting any number for it, ours included. Absent that sweep, the honest reference point is the closest apples-to-apples case NVIDIA has published: dense GEMM in CuTe DSL against dense GEMM in C++ CUTLASS, with no epilogue at all. NVIDIA's own benchmark found CuTe DSL matches C++ CUTLASS throughput on Blackwell across most shapes, with the K=512 exception described above, and trails C++ CUTLASS slightly on Ampere/A100. Since C++ CUTLASS GEMM is itself tuned to be competitive with cuBLAS on the shapes CUTLASS targets, that result is the best available signal for what a CuTe DSL versus cuBLAS comparison on a fused op should look like on Blackwell: close to parity on the matmul itself, with the epilogue fusion, which cuBLAS cannot do at all, as CuTe DSL's actual edge rather than raw GEMM throughput.

Run the baseline script above on your own shape and batch size, then add the epilogue to both sides, before deciding whether the fusion is worth the authoring cost for your workload; a 4096x4096 shape at BF16 is a reasonable starting sweep, but MLP epilogues in real models rarely land on a single clean square shape.

Reading the Protocol Against NVIDIA's and Lambda's Published Numbers

Two outside data points anchor this comparison. First, NVIDIA's GEMM parity result above is the ceiling case: it is dense GEMM with no epilogue, the simplest possible comparison. A fused epilogue kernel like the one in the protocol above should track close to that ceiling on Blackwell, since the epilogue add is cheap relative to the matmul itself, but expect the Ampere gap NVIDIA already reports to persist or widen slightly on A100, since CuTe DSL is behind C++ CUTLASS there before you even add a fused epilogue.

Second, FlashAttention-4's 2.7x margin over a Triton implementation of the same kernel is the outer bound of what warp-specialization on Blackwell can be worth, not a typical result. FA4 is an unusually well-optimized, purpose-built kernel from the authors of FlashAttention itself; a first custom GEMM epilogue you write is far more likely to land somewhere between "roughly matching cuBLAS with a free fusion" and "notably behind cuBLAS until you've iterated on tile sizes," which is exactly why profiling your own shape against both baselines, rather than trusting either published number directly, is the right next step before shipping.

Should You Write Custom Kernels or Just Rent More GPUs?

This is the question that actually decides whether any of the above is worth your team's time, and the honest answer for most teams is: not yet.

When Custom Kernel Work Pays Back

Custom CUTLASS or Triton kernel authoring pays back when three conditions hold together: a single operation dominates your GPU-hours (profiled, not guessed), the off-the-shelf library gap is measured and large, and the operation runs at enough volume that the engineering time amortizes. DeepGEMM is a good existing example of this bar being cleared: it is a CUTLASS-family FP8 grouped GEMM kernel purpose-built for MoE expert dispatch, and our DeepEP and DeepGEMM deployment guide covers the throughput gain it delivers over generic GEMM paths on H200 and B200 clusters at production MoE serving volume. FlashAttention-4 clears the same bar for attention. Both are cases where a small team wrote a kernel that now serves as infrastructure for the entire ecosystem, which is a very different cost equation than a one-off kernel for a single internal model.

When It Doesn't: Renting Your Way Past the Problem

For most inference and fine-tuning workloads, the faster and cheaper path is not writing a kernel, it is using the kernels already shipped inside the serving stack and renting the GPU generation that actually fits the job. FlashInfer already ships block-sparse attention, MLA, and FP4/FP8 quantized attention kernels as vLLM and SGLang's default backends on Blackwell, and TensorRT-LLM leans on cuBLAS and CUTLASS GEMMs internally for its engine builds, so a production deployment already gets CUTLASS-class performance without anyone on your team authoring a kernel. If a workload is memory-bound rather than compute-bound, the fix is very often a bigger GPU or a better batching configuration, not a custom kernel at all.

A Decision Checklist

Before starting custom kernel work, work through this in order:

  1. Profile first. Use Nsight Compute or the PyTorch Profiler to confirm which single operation is actually dominating GPU time. If nothing clears roughly 15-20% of total runtime, custom kernel work will not move your top-line cost.
  2. Check for an existing kernel first. FlashInfer, DeepGEMM, DeepEP, and the CUTLASS example repository already cover GEMM, grouped GEMM, and attention for the common cases. Writing a kernel that already exists is wasted effort.
  3. Measure the gap, don't estimate it. Benchmark the existing library against a naive PyTorch baseline on your actual shapes. A 10% gap rarely justifies kernel authoring; a 2-3x gap, the kind FA4 demonstrated against Triton, usually does.
  4. Pick the DSL by operation type. Memory-bound, pointwise, or reduction-heavy: Triton. Compute-bound GEMM, grouped GEMM, or attention: CuTe DSL, keeping the beta portability caveat in mind for anything you plan to keep running past the next CUTLASS release.
  5. If none of the above clears the bar, rent the right GPU instead. Testing a kernel idea across generations, or simply moving a memory-bound workload to a GPU with more HBM bandwidth, is often cheaper in engineering hours than the kernel authoring project itself.

That last option is worth pricing out concretely, since kernel-authoring time is expensive and GPU-hours are not. Bare-metal access also matters for the profiling step above: ncu's hardware performance counters and root-level nvidia-smi access need bare metal or --privileged containers, and a virtualized neighbor on a shared instance will distort a micro-benchmark enough to make a real 10% kernel gain look like noise.

As of this snapshot, Spheron's live catalog shows H100 80GB at $2.64/hr on-demand ($2.20/hr spot) and A100 80GB at $1.43/hr on-demand ($1.19/hr spot), both useful baselines for the profiling and correctness-validation phase before you commit to a Blackwell run. B300 288GB, the closest catalog match to the GB300 target CUTLASS 4.4's new CTK 13.1 examples were built for, lists at $10.27/hr on-demand ($5.85/hr spot); GB300 itself (Blackwell Ultra, SM103) is not a separate catalog line item, so confirm current availability before assuming B300 stands in for it exactly. B200 192GB is spot-only in this snapshot at $5.37/hr, worth checking again if you need guaranteed capacity for a longer benchmark campaign rather than a short validation run.

Pricing fluctuates based on GPU availability. The prices above are based on 04 Sep 2026 and may have changed. Check current GPU pricing → for live rates.


CuTe DSL is close to graduating out of beta, and on the operations it targets, GEMM and attention, it is already delivering C++-class throughput from a Python front end. Whether it belongs in your stack depends less on how NVIDIA describes it and more on whether you have actually profiled a compute-bound kernel that's worth the authoring time. For most inference workloads, the faster move is testing your model on the right GPU generation before writing a single line of CuTe DSL. See the Spheron quick-start guide for provisioning details.

A GEMM or attention kernel worth hand-authoring is worth benchmarking across GPU generations before you commit to one. Rent H100 SXM5 or B200 by the minute, with no contract, to run the comparison in this post yourself.

H100 on Spheron → | Check B200 availability → | Rent B300 →

FAQ / 04

Frequently Asked Questions

Not as a hard cutoff you can point to. NVIDIA's own CuTe DSL FAQ states the DSL will graduate out of public beta by the end of summer 2026, and during the beta it makes no portability promises between releases. As of this post's publish date, that window has effectively closed, so check the FAQ page directly for the current status before you pin a production build against it.

No. CuTe DSL is a Python-embedded domain-specific language, not a C++ API with Python bindings. You write tile-level kernel logic in Python, decorated with CuTe DSL's JIT, and it compiles down through the same MLIR-based pipeline as C++ CUTLASS. This is a common point of confusion because FlashAttention-4, the reference Blackwell attention kernel, is sometimes described as a 'C++ tile DSL' project when it is in fact written entirely in CuTe DSL's Python front end.

CuTe DSL supports NVIDIA architectures from Ampere (SM80) through Blackwell, including the Blackwell Ultra SM103 target (GB300) added in CUTLASS 4.4.0 under CUDA Toolkit 13.1. It does not add anything for older architectures like Volta or Turing.

Custom kernel work pays back when a single hot operation runs at massive volume and off-the-shelf libraries leave a large, measured gap versus the roofline. For most teams, especially early on, renting the right GPU generation and letting vLLM, SGLang, or TensorRT-LLM's built-in CUTLASS and FlashInfer kernels do the work is the faster path to a working, reasonably-priced system. Reach for custom kernel authoring only after profiling shows exactly where the gap is and what closing it is worth in dollars per token.

Try It Yourself

Try It on Real GPUs

The GPUs behind these guides are the ones you can rent here: H100s, H200s, B200s, and more, billed per minute with no contracts and no minimum. Pick one and you are live in under two minutes.

Deploy Time
< 2 min
Uptime SLA
99.9%
GPU Models
10+
Billing
Per-Min