A Case for Tracing Based DSL Kernel LanguagesGeorge's Blog
SearchSearch
Dark modeLight mode<br>Reader mode
!a.isFolder&&!b.isFolder||a.isFolder&&b.isFolder?a.displayName.localeCompare(b.displayName,void 0,{numeric:!0,sensitivity:\"base\"}):!a.isFolder&&b.isFolder?1:-1","filterFn":"node=>node.slugSegment!==\"tags\"","mapFn":"node=>node"}">Explorer
A Case for Tracing Based DSL Kernel Languages<br>May 26, 2026
On the architectural divide between parsing and tracing kernel DSLs, and what tends to go wrong in each.
The language for writing NVIDIA GPU kernels was always exclusively CUDA, but since Triton appeared, a wave of Pythonic DSLs has followed: CuTe-DSL, cuTile, Pallas, Gluon, Warp, and the more recent TileLang used in DeepSeek’s DeepGEMM. Most of these systems share the same goal of lowering a tile-oriented program into PTX or LLVM-IR, and are embedded in Python.
The question is how to embed the DSL into Python. Triton and CuTe-DSL parse the source AST. Pallas runs the function under abstract values and traces the resulting operations. (PyTorch’s torch.compile intercepts CPython bytecode rather than source, but that is still parsing, just against a smaller, post-desugared grammar; the same trade-offs apply.)
Most DSLs follow Triton’s lead and use parsing. This essay takes the alternative and argues that a tracing-based approach is often preferable.
CUDA and Templates
A CUDA kernel directly specifies the execution code for each thread. A textbook fused-softmax kernel in CUDA looks roughly like this:
template typename T, int BLOCK_SIZE><br>__global__ void softmax_kernel(const T* __restrict__ x,<br>T* __restrict__ y,<br>int n_cols) {<br>int row = blockIdx.x;<br>int tid = threadIdx.x;
__shared__ float sdata[BLOCK_SIZE];<br>const T* row_ptr = x + row * n_cols;
float local_max = -INFINITY;<br>for (int i = tid; i n_cols; i += BLOCK_SIZE)<br>local_max = fmaxf(local_max, float(row_ptr[i]));<br>sdata[tid] = local_max;<br>__syncthreads();<br>// ... tree reduction, exp, normalize, store ...<br>The element type T and the block size BLOCK_SIZE must be known at compile time, as __shared__ memory is statically sized, and the compiler must specialise loop bounds to enable vectorisation of the body. Hence any expansion of the supported configuration space multiplies the number of instantiations. Three element types and four block sizes already imply twelve instantiations, and the responsibility for dispatching among them rests with the caller.
Adding more templates and more generalisations to CUDA, one eventually reaches a heavily templated CUTLASS-like state.
CUTLASS: Building Blocks for CUDA Kernels
CUTLASS is what C++ template metaprogramming looks like when taken as a way to write GPU kernels. Consider the declaration of its principal Gemm class, the entry point most users first encounter, from include/cutlass/gemm/device/gemm.h:
template<br>/// Element type for A matrix operand<br>typename ElementA_,<br>/// Layout type for A matrix operand<br>typename LayoutA_,<br>/// Element type for B matrix operand<br>typename ElementB_,<br>/// Layout type for B matrix operand<br>typename LayoutB_,<br>/// Element type for C and D matrix operands<br>typename ElementC_,<br>/// Layout type for C and D matrix operands<br>typename LayoutC_,<br>/// Element type for internal accumulation<br>typename ElementAccumulator_ = ElementC_,<br>/// Operator class tag<br>typename OperatorClass_ = arch::OpClassSimt,<br>/// Tag indicating architecture to tune for<br>typename ArchTag_ = arch::Sm70,<br>/// Threadblock-level tile size (concept: GemmShape)<br>typename ThreadblockShape_ = typename DefaultGemmConfiguration<br>OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,<br>ElementAccumulator_>::ThreadblockShape,<br>/// Warp-level tile size (concept: GemmShape)<br>typename WarpShape_ = typename DefaultGemmConfiguration<br>OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,<br>ElementAccumulator_>::WarpShape,<br>// ... ten more parameters elided ...<br>bool ScatterD = false,<br>typename PermuteDLayout = layout::NoPermute><br>class Gemm { /* ... */ };<br>cutlass/gemm/device/gemm.h, lines 169–233. Around twenty template parameters, several with defaults that recursively look up DefaultGemmConfiguration.
A fragment of the canonical Hopper warp-specialized GEMM example shows how a user composes a kernel from nested CollectiveBuilders, each a template that pulls in dozens of further instantiations:
using namespace cute;
using TileShape = Shape_128,_128,_32>; // CTA tile<br>using ClusterShape = Shape_4,_2,_1>; // cluster of CTAs
using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<br>cutlass::arch::Sm90, cutlass::arch::OpClassTensorOp,<br>TileShape, ClusterShape,<br>cutlass::epilogue::collective::EpilogueTileAuto,<br>ElementAccumulator, ElementAccumulator,<br>ElementC, LayoutC, AlignmentC,<br>ElementC, LayoutC, AlignmentC,<br>cutlass::epilogue::collective::EpilogueScheduleAuto<br>>::CollectiveOp;
using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<br>ArchTag, OperatorClass,<br>ElementA, LayoutA, AlignmentA,<br>ElementB, LayoutB,...