NineToothed Arrange-and-Apply DSL

• NineToothed's Arrange-and-Apply DSL is a high-level programming model that enables developers to express complex tiled GPU kernels using serial Python code through a clear separation of compile-time arrangement and per-tile application.
• It utilizes tensor-oriented metaprogramming to symbolically manipulate tensor shapes and strides, converting meta-operations into optimized Triton kernels with negligible performance overhead.
• The DSL improves maintainability and reusability in ML kernel development by abstracting intricate tiling, broadcasting, and memory alignment strategies into modular, easy-to-read components.

NineToothed's Arrange-and-Apply domain-specific language (DSL) is a high-level programming model for machine learning compute kernels, designed to deliver serial programming semantics atop Triton's parallel code generation for GPUs. By structuring kernel code through a tensor-oriented metaprogramming (TOM) interface, NineToothed enables developers to describe complex, high-performance tiled computations in concise, maintainable serial Python while relying on automatic compilation to parallel primitives with minimal performance penalty. The paradigm is formalized around the explicit separation of "arrangement" (compile-time tensor meta-operation composition) and "application" (per-tile serial kernel logic).

1. Formal Model of Arrange-and-Apply

The formal foundation of Arrange-and-Apply in NineToothed is built on the symbolic manipulation of tensor shapes, strides, and meta-operations:

• Symbolic Tensors: For a $d$-dimensional tensor $T$, the model tracks $\operatorname{shape}(T) = (s_1,\dots,s_d)$ and $\operatorname{strides}(T) = (r_1,\dots,r_d)$, both as symbolic (compile-time) objects.
• Meta-Operations: The central meta-operation is $\mathit{tile}(T, t)$, with $t = (t_1, \dots, t_d)$ denoting tile sizes. Tiling produces a hierarchical tensor representation: an outer shape $o_i = \lceil s_i/t_i \rceil$ for $i=1 \dots d$ and a fixed inner tile shape $(t_1, \dots, t_d)$. This composition rewrites shape and stride attributes symbolically, constructing an abstract syntax tree (AST) of index expressions.
• Arrange-and-Apply Paradigm: Every compute kernel is decomposed into:
  • arrangement: Compile-time meta-op sequence yielding symbolic, hierarchically tiled tensors with matching outer shapes.
  • application: Serial, per-tile code (e.g., nested loops or linear algebra), treating arguments as concrete inner blocks.
  • Formally, for outer shape $\mathbf{o} \in \mathbb{Z}^m$, a grid of $|\mathbf{o}| = \prod_{i=1}^m o_i$ parallel kernel instances are launched, each invoking the application logic on respective tiles.

2. Front-End Syntax and Semantics

The front-end exposes a strictly two-phase user interface in pure Python, with explicit meta-programming constructs:

tensors = tuple(Tensor(d, name=f"t{i}") for i in range(k))

• Arrangement: Symbolic, meta-op-only Python functions. Example:

  def arrangement(A, B, C, BLOCK_M=block_size(), BLOCK_N=block_size()): C_arr = C.tile((BLOCK_M, BLOCK_N)) A_arr = A.tile((BLOCK_M, K)).expand((-1, C_arr.shape[1])) B_arr = B.tile((K, BLOCK_N)).expand((C_arr.shape[0], -1)) return A_arr, B_arr, C_arr

• Application: Per-tile serial code, written assuming input arguments are blocks (inner tiles):

  def application(A_blk, B_blk, C_blk): acc = ntl.zeros(C_blk.shape, dtype=ntl.float32) for k in range(A_blk.shape[0]): acc += ntl.dot(A_blk[k], B_blk[k]) C_blk = acc

• Kernel Construction and Invocation:

kernel = ninetoothed.make(arrangement, application, tensors) kernel(A, B, C, BLOCK_M=64, BLOCK_N=64)

This structure enables expressing complex tiling, broadcasting, and alignment strategies with concise serial code, leaving all tiling transformation to the compiler.

3. Compiler and Code Generation Pipeline

NineToothed employs a multi-stage transformation pipeline centered on AST manipulation and Triton IR generation:

• Frontend to Symbolic ASTs: Parses user "arrangement" and "application" code, constructing ASTs representing shape, stride, and tiling logic. No data is accessed; transformations are compile-time only.
• Tile-to-Program Mapping: Extracts the outer tiling shape $\mathbf{o}$ and emits a Triton kernel using grid-based launching:

  @triton.jit def kernel(...): pid = triton.program_id(axis=0) idx = unravel(pid, o)

  For each tensor, records dimension mappings ("source_dims", "target_dims") for correct block alignment and broadcasting.
• Source-to-Target Mapping and Codegen: For each parameter:
  • Computes per-tile memory offsets:

  $$\mathbf{off} = \sum_{i=1}^{d} \bigl(\operatorname{idx}[i] \cdot t_i + \operatorname{arange}(t_i)\bigr)\times \operatorname{stride}_i$$ - Emits triton.load/triton.store operations to wrap application logic. - Inlines the user application's serial code directly into the kernel. Final Triton ASTs are lowered to optimized GPU code via the Triton LLVM/MLIR toolchain.

4. Transformations and Algorithmic Workflow

The translation from high-level Python to GPU code follows an explicit algorithmic structure:

Step	Workflow Component	Description
1	Symbolization	Instantiate symbolic Tensors for each kernel input
2	Arrangement	Call user-defined arrangement to build arranged tensors
3	Outer Shape Agreement	Assert all arranged tensors share the same outer tiling
4	Kernel Launch	Emit Triton kernel mapped over the outer tile grid
5	Block Offsets	For each argument, compute block-local offsets for loads/stores
6	Application Inlining	Inline serial application logic using Triton AST nodes
7	Output	Generate stores and finalize the Triton AST

This transformation underlies the "arrange" $\to$ "apply" regime, providing a modular and composable code path from symbolic meta-operations to efficient GPU kernel implementations (Huang et al., 16 Jul 2025).

5. Performance Characteristics

Quantitative results on an NVIDIA A100 GPU demonstrate the impact of the Arrange-and-Apply DSL on maintainability, code size, and performance:

• Code Metrics:

  • Lines of meaningful code (LLOC) reduced to 0.25%–56% of analog Triton kernels.
  • Halstead volume reductions reach ~96%.
  • Maintainability Index improved across all tested kernels (e.g., 44 to 60+).
• Kernel Microbenchmarks:
  • Max deviations from hand-crafted Triton across all shapes are +3.93% / –1.58%.
  • Average difference is ~0.4%.
• End-to-End ML Inference:
  • In a Transformer with Attention, Linear, RMSNorm, and SiLU layers replaced, total throughput impact ranges from –5.3% to +0.3% (average –1.8%).

This suggests NineToothed's abstractions impose negligible runtime overhead and attain parity with expert-optimized Triton code (Huang et al., 16 Jul 2025).

6. Advantages and Constraints

Advantages

• Serial Programming Model: Enables writing GPU kernels without explicit parallel programming constructs (program_id, pointer math, warp index manipulation).
• Conciseness and Readability: Meta-operations like tile, expand, and flatten allow complex memory access and tiling to be expressed succinctly and safely.
• Reusability and Modularity: The formal separation of arrangement and application supports kernel logic reuse (e.g., leveraging GEMM arrangements in multiple contexts such as convolutions).
• Maintainability: Reduced code complexity and improved maintainability indices.
• Performance Parity: Empirically matches Triton hand-written kernel performance on standard ML workloads.

Limitations

• Compile-Time Overhead: Construction and symbolic analysis of ASTs in Python introduce extra latency relative to writing raw Triton code directly.
• Reduced Low-Level Control: Certain advanced optimizations, unorthodox tiling, or use of vendor-specific instructions are not directly accessible.
• Limited Meta-Op Suite: Only a fixed set of meta-operations are supported; kernels with extremely irregular access patterns may require fallback to custom Triton code.
• Reliance on Triton Tuning: Kernel efficiency depends on auto-tuning heuristics and transformations provided by the underlying Triton backend.

A plausible implication is that NineToothed is especially well-suited for rapid, high-level prototyping of dense tensor kernels, with hand-coded Triton or lower-level languages preferable in cases demanding non-standard optimizations or full control of mapping details.

7. Context and Significance

The introduction of the Arrange-and-Apply paradigm in NineToothed addresses persistent challenges in kernel development for machine learning systems: bridging high-level, serial tensor computation expressions with the requirement for high-performance, parallel execution. By shifting from hand-authored parallelization and memory scheduling to a meta-op-driven code generation flow, the approach advances the maintainability, productivity, and reliability of ML kernel programming. The negligible end-to-end performance cost demonstrated across representative workloads positions Arrange-and-Apply as a generative model for kernel DSLs in the deep learning ecosystem, evolving the role of domain-specific languages such as Triton towards greater accessibility and code quality (Huang et al., 16 Jul 2025).