PaperlayerView exact revision on arXiv

Source-linked AI summary

TileBench: A Controlled Benchmark for Performance Evaluation and Bottleneck Diagnosis of Tile-Based Programming Models

Bowen Cui, Zhongchun Zhou, Hao Wu, Tejas Ramesh, Junyu Yin, Jialiang Gu, Keren Zhou

arXiv:2609.29067v1cs.PFcs.PL

TL;DR

Tile-based GPU programming models are difficult to compare systematically in performance, tuning, and usability across AI workloads. TileBench provides a controlled comparison of Triton and cuTile, finding workload-dependent performance strengths and greater token efficiency for Triton in iterative LLM kernel generation.

  • Problem

    Performance bottlenecks and comparisons between tile-based DSLs remain difficult to understand systematically across diverse AI operators.

    Source: §1 · p. 1

  • Method

    TileBench evaluates 45 manually written, verified Triton and cuTile implementations under matched semantics and standardized dtype and input-size sweeps.

    Source: §1 · p. 2

  • Results

    Triton is more robust on irregular, streaming, and bandwidth-bound operators, while cuTile excels on workloads matching static tiles, Tensor Cores, and reusable tile movement; Triton is also more token-efficient in LLM generation.

    Source: §1 · p. 2, §6 · p. 9

  • Takeaways & Limitations

    Neither programming model is uniformly superior: their relative performance varies with workload characteristics.

    Source: §6 · p. 9

    TileBench currently covers one NVIDIA B200 GPU and two tile-based DSLs, and its analytical roofline metrics do not replace full NCU-counter analysis.

    Source: Limitations · pp. 9–10

Abstract

from arXiv · show

Tile-based programming models, such as Triton and cuTile, aim to simplify high-performance kernel development, but their practical performance, tuning behavior, and usability remain difficult to compare systematically. We present TileBench, a controlled benchmark for evaluating Triton and cuTile on NVIDIA B200 GPUs under matched operator semantics and comparable implementation structures. TileBench contains 45 operators covering diverse AI-kernel patterns and memory/computation behaviors. Each task provides a PyTorch reference, verified Triton and cuTile implementations, standardized data-types (dtype) and input-size sweeps, default and autotuned configurations, roofline-based metrics, and profiling-guided diagnosis. Our evaluation shows that performance gaps are workload-dependent: cuTile excels on a small cluster of Tensor-Core/TMA-friendly kernels, while Triton is stronger on many irregular, streaming, and bandwidth-bound operators. We further evaluate LLM-generated cuTile and Triton kernels and find that Triton is consistently more token-efficient than cuTile under the same iterative refinement protocol. TileBench is publicly available at https://github.com/Deep-Learning-Profiling-Tools/Tilebench.

1 Introduction

TileBench addresses the difficulty of diagnosing and comparing tile-based GPU kernels across diverse AI operators. It evaluates Triton and cuTile through matched benchmark tasks and reports workload-dependent performance differences, including a token-efficiency advantage for Triton in LLM-generated kernels.

Benchmark contribution

  • TileBench provides 45 operators with manually written, verified Triton and cuTile implementations under matched semantics and comparable implementation structures.The benchmark uses standardized dtype and input-size sweeps on an NVIDIA B200 GPU.Source: §1 · p. 2

Evaluation and diagnosis

  • TileBench evaluates default and autotuned configurations to measure tuning gains and remaining headroom relative to hardware roofline limits.A profiling-guided workflow uses generated code, compiler behavior, and hardware metrics to investigate performance differences.Source: §1 · p. 2

LLM-generated kernels

  • The LLM-generated-kernel track measures correctness, runtime performance, and the token consumption needed to reach performant implementations.It applies the same iterative refinement protocol to Triton and cuTile.Source: §1 · p. 2

Main findings

  • Across the suite, both DSLs substantially outperform PyTorch, Triton is stronger overall, and cuTile is competitive on a small set of Tensor Core- and TMA-friendly workloads.Autotuning improves both backends but leaves substantial roofline headroom; Triton is also consistently more token-efficient than cuTile in the iterative LLM-generated-kernel evaluation.Source: §1 · p. 2

2 Related Work

TileBench focuses on Triton and cuTile because both expose tile-centric interfaces and can be compared on the same NVIDIA B200 platform. It complements existing kernel-generation benchmarks by isolating programming-model effects between tile-level DSLs.

2.1 Tile-Based GPU Programming Models

  • Triton and cuTile both expose tile-centric programming interfaces for NVIDIA GPUs, but differ in their APIs and compiler/runtime mappings to GPU operations.Triton exposes a Python-embedded block program; cuTile expresses kernels through tile operations such as ct.load, ct.mma, and ct.store.Source: §2.1 · p. 2
  • TileBench focuses on Triton and cuTile because both run on B200 and support matched-semantics comparison with comparable tuning budgets, roofline analysis, and profiling-guided diagnosis.Other systems are not all directly comparable in this setting because of different accelerator targets or lower-level control.Source: §2.1 · pp. 2–3

2.2 Benchmarks for GPU Kernel Generation and Optimization

  • Existing GPU-kernel benchmarks measure generation or optimization but do not isolate programming-model effects between tile-level DSLs.TileBench's main track compares manually written, verified implementations, while a separate track measures LLM generation ease and token efficiency.Source: §2.2 · p. 3

3 TileBench

TileBench standardizes comparison of Triton and cuTile through a shared task format, evaluation harness, benchmark suite, and LLM-generation track. It reports performance and roofline metrics and uses profiling for selected performance gaps.

3 TileBench

  • TileBench compares Triton and cuTile under matched operator semantics and standardized evaluation, using PyTorch as the semantic reference, correctness oracle, and software baseline.The tile-based programming models are the primary comparison targets.Source: §3 · p. 3

3.1 Overview and Pipeline

  • The workflow curates operators into a unified task format, runs PyTorch, Triton, and cuTile through a shared execution harness, then aggregates comparative metrics.For selected large-gap cases, Nsight Compute and PTX/SASS inspection examine instruction mix, memory behavior, TMA/tcgen05 usage, and bank conflicts.Source: §3.1 · p. 3

3.2 Benchmark Suite

  • The 45-task suite standardizes implementations, dtypes, input shapes, correctness checks, timing, and comparable autotuning across Triton and cuTile.Tasks originate from TritonBench and LeetGPU; each operator receives 20 input configurations per supported dtype, and the backends use comparable—not identical—tuning ranges.Source: §3.2 · pp. 3–4

3.3 Iterative LLM Code-Generation

  • The LLM track generates each backend's implementation from task specifications, refines it with verification and performance feedback, and selects its best verify-clean iteration.Each backend runs up to 10 iterations on a representative large case per supported dtype, without backend autotuning; refinement freezes after verification and at least 80% roofline utilization or budget exhaustion.Source: §3.3 · pp. 4–5

3.4 Evaluation Metrics

  • The evaluation reports latency, PyTorch speedup, throughput, effective bandwidth, arithmetic intensity, and roofline utilization; token efficiency is correctness-gated speedup per million tokens.Non-compiling or incorrect iterations contribute zero speedup while still incurring token costs; roofline bounds use dtype-specific B200 peak throughput and peak HBM bandwidth.Source: §3.4 · p. 5

4 Experimental Setup

The evaluation measures Triton and cuTile on one NVIDIA B200 GPU and uses consistent formulas and aggregation for benchmark metrics. A separate LLM track evaluates two reasoning models.

Evaluation tracks

  • The evaluation has a main track for human-written Triton and cuTile implementations across all 45 operators and a separate iterative LLM code-generation track.The LLM track uses GPT-5.5 and Claude Opus 4.7 through their respective APIs with default decoding settings.Source: §4 · pp. 5–6

Hardware and software

  • All measurements use one NVIDIA B200 GPU with 180 GB of HBM3e memory, and roofline analysis uses a measured peak bandwidth of 6539.4 GB/s.The software stack is PyTorch 2.10, CUDA 13.0, Triton 3.6.0, and cuda-tile 1.3.0.Source: §4 · p. 6

5 Evaluations •

The evaluation compares autotuned Triton and cuTile against PyTorch, examines tuning gains and remaining hardware headroom, diagnoses performance gaps, and tests LLM-generated kernels under a matched refinement protocol. Across these questions, results vary by workload, while Triton is stronger overall and more token-efficient in the LLM track.

5.1 RQ1: Autotuned Programming-Model Performance

  • Triton achieves a 2.02× geometric-mean speedup over PyTorch versus cuTile’s 1.58×, and is faster than cuTile on 37/45 operators.CuTile is faster on 8/45 operators; its clearest advantages occur on regular, reuse-heavy tiled workloads, including Tensor-Core/TMA kernels.Source: Figure 2 · p. 6, §5.1 · pp. 6–7

5.2 RQ2: Autotuning Sensitivity and Performance Headroom

  • Autotuning improves roofline utilization by 1.18× for Triton and 1.22× for cuTile, but substantial hardware headroom remains.Only 7/45 Triton and 5/45 cuTile implementations reach at least 80% operator-level geometric-mean roofline utilization; tuning covers declared configuration spaces, not alternative algorithms or lowering strategies.Source: Figure 3 · p. 7, §5.2 · p. 7

5.1 RQ1: Autotuned Programming-Model Performance

  • PyTorch speedups do not isolate compiler-backend effects because eager fusion, shape specialization, and vendor-library dispatch also affect the comparison.PyTorch can use optimized implementations such as cuDNN convolution, cuBLAS GEMM, CUB radix sort, and warp-specialized attention kernels.Source: §5.1 · p. 7

5.3 RQ3: Performance-Gap Diagnosis

  • All 20 largest sweep-max latency gaps favor Triton, while cuTile is 1.55× faster on FP32 matmul, with gaps tied to indexing, reduction, state-placement, and scheduling differences.The largest Triton-faster gaps include flash_decode/FP32 at 3.72× and block_sparse_attention/FP16 at 3.08×; cuTile’s clearest advantages arise on regular tiled computation.Source: §5.3 · pp. 7–9, Figure 4 · p. 8
  • In weight_dequant and destindex, cuTile’s defensive index arithmetic drives instruction bloat, while destindex also shows less-coalesced loads than Triton.CuTile executes 6.2× more instructions in BF16 weight_dequant and 17.3× more in Int8 destindex; destindex has 1.78 sectors per load request for cuTile versus 8.50 for Triton.Source: §5.3 · p. 8

5.4 RQ4: LLM Usability and Token Efficiency

  • Under the maximum 10-iteration protocol, Triton yields more verify-clean kernels, higher correctness-gated speedups, lower token costs, and higher token efficiency for both evaluated LLMs.Triton’s token-efficiency advantage over cuTile is approximately 47% for GPT-5.5 and 70% for Claude Opus 4.7.Source: §5.4 · p. 9

6 Conclusions

TileBench is a reproducible benchmark for comparing tile-based GPU programming models. Its results show workload-dependent strengths for Triton and cuTile, alongside tuning and LLM-generation findings.

6 Conclusions

  • TileBench is presented as a reproducible benchmark for comparing tile-based GPU programming models and studying tile-centric DSLs.Source: §6 · p. 9
  • Both Triton and cuTile substantially outperform PyTorch, but Triton is stronger on irregular, streaming, and bandwidth-bound operators while cuTile excels on Tensor-Core and reusable-tile-movement workloads.Neither abstraction is uniformly superior.Source: §6 · p. 9
  • Autotuning yields measurable gains, yet most kernels remain far from roofline limits, with backend lowering and memory or instruction behavior also mattering.The passage highlights memory staging, instruction selection, and shared-memory behavior as factors beyond tile-size search.Source: §6 · p. 9
  • The LLM-generated track finds that programming-model design affects final performance and the token cost of reaching correct, efficient kernels.Source: §6 · p. 9

Limitations

TileBench’s current scope is limited to one GPU and two tile-based DSLs. Its roofline metrics are standardized analytical summaries, not substitutes for full NCU-counter analysis.

Limitations

  • TileBench currently targets a single NVIDIA B200 GPU and two tile-based DSLs.Source: Limitations · p. 9
  • Its roofline metrics use analytical FLOP/HBM-byte formulas.Source: Limitations · p. 10
  • Those metrics provide a standardized hardware-aware summary rather than replacing full NCU-counter analysis.Source: Limitations · p. 10

A.1 Full Benchmark List

This appendix documents TileBench’s benchmark inventory, configuration and B200 assumptions, and the materials and requirements for iterative LLM kernel generation, including API references and operator problem descriptions.

A.2 Operator Category

  • Per-operator verification tolerances reflect differing shape-sensitivity profiles and may be widened for floating-point reassociation in non-associative reductions.The setup exercises matmuls where tensor-core paths activate, reductions near occupancy saturation, and elementwise operations across memory- and compute-bound regimes.Source: §A.2 · p. 12, §A.3 · p. 13

A.3 Config.yaml Example

  • Each operator’s config.yaml declares the harness case grid as the cross product of input sizes and dtypes.The operator categories are documented in Table 3.Source: §A.3 · p. 12, Table 3 · p. 13
  • TileBench contains 45 operators: 26 adapted from TritonBench and 19 from LeetGPU.Table 1 and Table 2 identify the operators derived from each source.Source: Table 1 · p. 12, Table 2 · p. 12, Table 3 · p. 13

A.4 B200 JSON Information

  • The B200 peak bandwidth baseline is 6539.4 GB/s, measured on the system; FP32 dot/MMA workloads use the dense TF32 Tensor Core ceiling.Dense compute ceilings are derived from NVIDIA HGX B200 / DGX B200 specifications and exclude 2:4 sparsity.Source: §A.4 · p. 13

B.1 Iteration 0 Prompt Template

  • The iteration-0 prompt concatenates framework conventions, backend API references, the operator description, config.yaml, the PyTorch reference, and output-format instructions, with no prior model context.Operator-specific material is supplied in the user message; vector_add is the running example for the prompt documentation.Source: §B.1 · p. 13, B Iterative LLM Code-Generation Prompts · p. 13

B.2 Iteration N Prompt Template (N > 0)

  • For later iterations, prompts are rebuilt with the trajectory, previous source files, and—after regression or verification failure—the best verify-clean implementation, followed by the shared task materials.The task materials include the framework guide, API references, operator description, config.yaml, PyTorch reference, and output-format instructions.Source: §B.2 · p. 14

B.3 TileBench Framework Guide

  • Generated implementations follow the run() / get_last_config() contract and use one hard-coded configuration per backend, with up to 10 refinement iterations driven by prior performance feedback.A backend freezes after verification passes and reaches at least 80% roofline utilization; the best verification-passing iteration becomes the final result.Source: §B.3 · pp. 15–16

B.4 Triton and cuTile API References

  • The API references scope generated code to Triton 3.6.0 and TileBench’s cuTile packages and cover programming, memory, compute, synchronization, constraints, and the harness contract.The references differ in their topical sections to reflect the distinct abstractions exposed by Triton and cuTile.Source: §B.4 · p. 16
  • The cuTile memory-operations reference presents function signatures, parameters, and representative calls, including load/store behavior and options for ordering, padding, latency hints, and TMA.Out-of-bounds stores are silently ignored.Source: §B.4 · p. 17

B.5 Operator Problem Description

  • Each operator’s problem description specifies its mathematical definition, input and output contracts, and a worked example; vector_add defines element-wise addition of equal-length 1D arrays.The vector_add example also requires native features, an unchanged solve signature, and storing the result in output.Source: §B.5 · pp. 17–18

B.7 Output Format Instructions

  • Generated implementations must use the prescribed fenced Python file format for each active backend, with no extra code blocks or prose beyond a brief approach summary.The output instructions exclude test code and PyTorch reference code.Source: §B.7 · p. 18

B.6 PyTorch Reference Implementation

The PyTorch reference specifies each operator’s input/output contract and serves as the correctness oracle. For vector_add, the reference returns x + y; the harness checks generated-kernel signatures and output tolerances.

B.6 PyTorch Reference Implementation

  • The PyTorch reference serves as both the operator input/output specification and the correctness oracle for generated kernels.Source: §B.6 · p. 18
  • For vector_add, the reference implementation returns x + y.Source: §B.6 · p. 18
  • The harness checks that generated run() functions match the reference’s positional signature and output within config.yaml verification tolerances or per-dtype defaults.Source: §B.6 · p. 18

C.1 Evaluation Modes and Aggregation Conventions

The protocol separates per-case timing from geometric-mean aggregates, then uses profiling evidence to diagnose performance differences while retaining PyTorch as the practical software baseline.

C.1 Evaluation Modes and Aggregation Conventions

  • Each operator–backend–dtype–case latency averages 100 timed launches after 20 warmups, with CUDA graph replay and L2 eviction before each launch.The complete operator run is timed as one unit, preserving intended producer-to-consumer cache reuse within multi-kernel implementations.Source: §C.1 · p. 18
  • Paper-level aggregates use geometric means, weighting valid dtype–case pairs equally per operator and all 45 operators equally in the suite.Category results use the same construction over operators in each category; medians are distribution statistics, not averages.Source: §C.1 · pp. 18–19

C.2 NCU Profiling Protocol and Kernel Attribution

  • NCU profiles the maximum catalogued case for each supported operator–dtype pair and diagnoses gaps by relating source, generated code, hardware counters, and Proton latency.NCU counters explain execution mechanisms; benchmark latency remains based on Proton measurements, not NCU kernel duration.Source: §C.1 · p. 19, §C.2 · p. 19

C.1 Evaluation Modes and Aggregation Conventions

  • RQ4 reports correctness-gated, single-case speedup, best verify-clean speedup, tokens consumed, and best-speedup-per-million-token efficiency.Skipped iterations after a backend freezes add neither tokens nor speedup.Source: §C.1 · p. 19

C.2 NCU Profiling Protocol and Kernel Attribution

  • Diagnosis is limited to the profiled implementations, compiler versions, tuned winners, and B200; dtype, tile shape, layout, and compiler choices can change instruction paths.Higher-level compiler IRs were unavailable in finalized artifacts, so lowering claims rely on PTX, SASS, and source-correlated evidence; source-family counts may be lower bounds.Source: §C.2 · p. 20

C.3 PyTorch Baseline

  • PyTorch provides the semantic reference, correctness oracle, and practical software baseline, whose execution path varies by operator and may use vendor libraries or other kernels.Speedup over PyTorch reflects the custom kernel’s benefit over that available path; RQ3 primarily compares matched Triton and cuTile implementations.Source: §C.3 · p. 20

D.1 RQ1: Autotuned Speedup Distributions and Category Breakdowns

Triton leads category-level speedup across all five groups, though cuTile wins several regular tiled cases. Autotuning yields modest suite-level gains, while substantial backend headroom remains for much of the suite.

D.1 RQ1: Autotuned Speedup Distributions and Category Breakdowns

  • Point-wise operators have the highest category-level roofline utilization and many near-parity results, while data-layout operators have large PyTorch speedups but much lower utilization.Stencil/convolution and matrix-multiplication/attention are least utilized; their limits include library gaps, dynamic indexing, state placement, and pipeline formation.Source: §D.1 · p. 20
  • 1.95× vs 1.58× median operator-level speedup for Triton and cuTile, respectively; both reach approximately 36× on bitonic_sort.That bitonic_sort speedup partly reflects the eager multi-pass PyTorch reference rather than only DSL backend efficiency.Source: §D.1 · p. 20, Figure 8 · p. 21
  • Triton has the highest category-level speedup in all five groups, while cuTile wins regular tiled cases including matmul_fp32_fp16_fp8, flash_attention, and jacobi_stencil_2d.CuTile’s category aggregates are reduced by deficits on irregular indexing, small reductions, sparse traversal, and runtime-controlled loops.Source: §D.1 · p. 20, Table 4 · p. 21

D.2 RQ2: Per-operator Autotuning Gains and Roofline Headroom

  • 1.18× vs 1.22× suite-level geometric-mean tuning gain for Triton and cuTile, respectively, but median gains are 1.07× and 1.06×.The larger cuTile geometric mean comes from several high-gain operators rather than a uniformly larger shift.Source: Figure 9 · p. 21, §D.2 · p. 21
  • Four Triton and four cuTile operator results fall below their defaults after tuning; seven of these eight regressions are within approximately 1.5% of parity.Triton top_k_selection is the main exception, at 0.92× its default.Source: §D.2 · p. 22
  • Configuration headroom comes from tuning parameters, but backend headroom persists when source and compiler choices produce costly instruction paths or resource use.The small number of operators above 80% utilization after tuning indicates that backend headroom dominates much of the suite.Source: §D.2 · p. 22

D.3 RQ4: Iterative Refinement Trajectories

Under the 10-iteration protocol, generated Triton kernels achieve verify-clean coverage on all 45 operators, while cuTile coverage is lower. The section also reports the experiments’ compute and API-token budgets and artifact licensing boundaries.

D.3 RQ4: Iterative Refinement Trajectories

  • Later refinement improved some trajectories, but the fixed 10-iteration budget cannot establish whether more iterations would close the backend gap.The evaluation also does not isolate API expressibility, primitive availability, compiler constraints, documentation, examples, or prior model exposure as causes of the usability difference.Source: Figure 10 · p. 22, §D.3 · p. 22

E Computational Budget

F LLM API Usage

  • GPT-5.5 consumed 18.68M API tokens at $279.05, while Claude Opus 4.7 consumed 21.09M tokens at $191.66.Together, the LLM track used 39.77M tokens and cost $470.71; both evaluated models are closed-source API models with unavailable parameter counts.Source: F LLM API Usage · p. 22

G Artifact Sources and Licenses

  • TileBench preserves attribution and license notices for TritonBench-derived components and independently writes LeetGPU-derived task artifacts rather than redistributing LeetGPU source or tests.LeetGPU’s CC BY-NC-ND 4.0 license permits non-commercial sharing with attribution but not redistribution of adapted material.Source: G Artifact Sources and Licenses · pp. 22–23

D.3 RQ4: Iterative Refinement Trajectories

  • Triton achieved verify-clean implementations for all 45 operators with both models; cuTile achieved 43/45 with GPT-5.5 and 41/45 with Claude Opus 4.7.For the remaining cuTile operators, none of the generated implementations within the 10-iteration budget passed verification.Source: §D.3 · p. 22, Table 6 · p. 23
Loading 2609.29067v1…