Correct but Slow: An Empirical Study of the GPU Kernel Evaluation Gap in Modern Domain-Specific Languages

Abstract: Modern GPU domain-specific languages (DSLs), such as Triton and TileLang, are increasingly used to implement specialized deep-learning kernels and as target languages for automated kernel-generation systems. Existing DSL-kernel evaluations establish correctness through reference-based numerical validation -- necessary, but silent on replacement quality: a functionally valid kernel may still fall far below the throughput of the optimized library operator it is intended to replace. We study this correctness-performance gap using 22 Triton and TileLang kernels from five operator categories on NVIDIA A100 and GH200 GPUs, asking whether correctness-based evaluation identifies kernels unsuitable as library replacements, why such failures occur, and how they can be detected without exhaustive benchmark coverage. The study yields three results. \emph{First}, correctness-based evaluation can admit severe slowdowns: an idiomatic TileLang LayerNorm kernel passes KernelBench's correctness check while running more than 300$\times$ slower than the PyTorch baseline. \emph{Second}, the causes differ by kernel family. TileLang normalization and reduction slowdowns are mainly repairable authoring defects, such as sequential reductions and unnecessary dtype conversions, whereas convolution and large general matrix multiplication (GEMM) retain residual gaps after optimization due to code-generation and autotuning-coverage limits; vendor-library algorithm selection contributes only marginally. \emph{Third}, two lightweight checks -- library-relative efficiency and roofline utilization -- are complementary screening criteria: together they flag every functionally valid but inefficient kernel in our suite and separate repairable authoring defects from structural residuals.

• The paper demonstrates that correctness-centric benchmarks mask severe performance inefficiencies, with some DSL kernels running up to 1520× slower than vendor standards.
• Empirical evaluations on 22 kernels using NVIDIA A100 and GH200 show that authoring defects, absent vectorization, and autotuning gaps are key contributors to performance issues.
• The study introduces lightweight comparability and roofline-based heuristics that guide optimizations and enhance the performance reliability of GPU kernels.

Correct but Slow: An Empirical Analysis of GPU Kernel Evaluation Gaps in Modern DSLs

Introduction and Context

The paper "Correct but Slow: An Empirical Study of the GPU Kernel Evaluation Gap in Modern Domain-Specific Languages" (2607.04454) provides an exhaustive empirical investigation into the longstanding disconnect between functional correctness and performance efficacy in GPU kernel evaluation—specifically within emerging domain-specific languages (DSLs) such as Triton and TileLang. As deep learning frameworks demand custom, model-specific, and fused GPU kernels that go beyond the static operator sets offered by vendor libraries (e.g., cuBLAS, cuDNN), these DSLs play an expanding role both for expert developers and as output targets for LLM-based kernel generation workflows.

While correctness-centric benchmarks (notably KernelBench and TritonBench) have become de facto arbiters of kernel validation, their reliance on reference-based numerical checks obscures a critical issue: a kernel may pass all functional requirements yet perform orders of magnitude slower than its vendor-library counterpart or even fail to leverage available hardware resources. This paper rigorously delineates the dimensions of this "correct-but-slow" phenomenon, investigates the roots of admitted performance gaps, and develops lightweight, actionable heuristics for practitioners targeting performant and portable GPU code with DSLs.

Methodology and Evaluation Design

To structurally probe the correctness–performance gap, the study selects 22 kernels spanning five operator categories (GEMM, attention, convolution, normalization, element-wise/reduction) and implements both DSL (Triton, TileLang) and vendor-library baselines on NVIDIA A100 and GH200 GPUs. Beyond latency metrics, the analysis incorporates detailed hardware profiling (using Nsight Compute counters) to attribute performance deviations to root causes—including authoring mistakes, code-gen limitations, and fundamental library-maturity boundaries. Notably, all considered kernels pass strict reference-based numerical validation, ruling out functional defects as an explanation for poor performance.

A representative example underscores the implications: an idiomatic TileLang LayerNorm kernel passes all correctness checks but is measured at over 300× slower than the PyTorch (cuDNN-based) baseline. The study explores not only the magnitude of admitted slowdowns but also their provenance across operator families. Optimization campaigns, including LLM-based agentic kernel tuning, are leveraged to validate the recoverability of observed gaps.

Empirical Findings: Existence and Magnitude of the Evaluation Gap

The core quantitative finding is that correctness-gated benchmarks do not distinguish performant kernels from functionally valid but highly inefficient implementations. This gap is both architecturally and operator-dependent:

• Normalization and Reduction: TileLang kernels, in default idiomatic implementations, see catastrophic collapses—LayerNorm and RMSNorm exhibit slowdowns up to 1347× and 1520×, respectively, compared to baseline (PyTorch). Contrastingly, Triton kernels are at or near parity for the same tasks.
• Convolution and Large GEMM: Both DSLs underperform vendor libraries significantly, though TileLang generally fares better than Triton for convolution. Performance losses on large-square GEMM are primarily due to absent autotuning coverage and cache utilization differences.
• Element-wise Operators: Most kernels are within a reasonable margin (60–98%) of baseline, with outliers arising from index-returning reductions (e.g., argmax) and certain fusion patterns able to exceed library performance due to more aggressive specialization in the DSL code.

Figure 1: Library efficiency (%) of Triton and TileLang kernels vs. cuBLAS/cuDNN on the A100, by category (log scale; dashed line = 100% parity).

The pattern reappears across architectures, as confirmed by profiles on the NVIDIA GH200:

Figure 2: Library efficiency (%) of Triton and TileLang kernels relative to cuBLAS/cuDNN on the NVIDIA GH200-480GB, by category (log scale; dashed line = 100% parity).

A critical observation is that the performance gap is not uniform. For example, while Triton normalization kernels achieve 90% of baseline, normalization and reduction in TileLang can fall below 1%. For convolution, both DSLs are weaker than cuDNN, with gaps widening at larger filter sizes and for depthwise modes due to launch overhead and absent codegen vectorization.

Root Cause Analysis

The paper systematically traces observed performance gaps to a small taxonomy of root causes:

• RC0 (Authoring Defects in TileLang Reductions/Norms):
  • Failure to convert sequential reduction loops (T.serial) to parallel reductions (T.reduce) results in massive exposed memory latency, as evidenced by HTML-reported hardware counters (long scoreboard stalls).
  • Compiler-generated non-vectorized loads exacerbating memory traffic.
• RC1 (Codegen: Absent Vectorization of Strided Accesses):
  • Convolution kernels lack wide/contiguous memory accesses, precluding use of cp.async and collapse into scalar loads, severely underutilizing bandwidth.
• RC2 (Autotuning and Caching Inadequacies):
  • Triton's auto-tuner omits configurations critical at large shapes or for L2 cache reuse, leaving kernels bandwidth-bound or unable to exploit hardware efficiently.
• RC3 (Register Pressure and Occupancy):
  • Certain kernels (TileLang LayerNorm) experience register spill collapse, yielding sub-20% occupancy due to compiler allocation choices.
• RC4 (Library Algorithm Diversity):
  • Vendor libraries benefit from algorithmic diversity (Winograd for convolution), which is missing in DSL kernels but, empirically, accounts for only a small percentage of the residual gap.

Each root cause leaves a distinctive counter signature, providing practitioners with actionable debugging pathways (e.g., spikes in warp_stall_long_scoreboard pinpoint memory-latency bottlenecks amenable to reduction parallelization).

Heuristics for Performance Validation and Repair

Given that exhaustive benchmarking is infeasible across the combinatorial (DSL × data type × shape × GPU) space, the authors introduce two practical heuristics to flag and guide repair of performance-poor kernels:

• Comparability Screen (Baseline-Relative):
  • Rapidly flags catastrophic slowdowns by comparing a candidate kernel’s performance against the vendor-library reference over representative input shapes.
  • Limitations: this is not a certificate; it may over-credit kernels outperforming a non-fused baseline.
• Roofline Anchor (Hardware-Peak-Relative):
  • Computes the achieved work rate as a fraction of theoretical peak (using the roofline model [williams2009roofline]), providing a baseline-independent measure of kernel efficiency.
  • Crucially, these heuristics are shown to be complementary, identifying all functionally correct yet grossly underperformant kernels in the suite.

Optimization passes driven by these heuristics recover almost all dramatic collapses to within ±5% of library parity for normalization and reduction, by applying a small set of recurring patterns (e.g., conversion from serial to parallel reductions, avoidance of unnecessary datatype casting). For other gap categories—large GEMM and convolution—the residuals persist after best-effort optimization, underscoring genuine maturity and autotuning coverage differences between vendor libraries and DSL infrastructure.

Implications and Prospective Developments

For DSL and Compiler Developers

• Enhanced code generation for vectorized memory access, especially in strided or spatially indexed patterns, is critical for closing the convolution gap.
• Auto-tuning mechanisms should be expanded to better cover high-utilization configurations, especially with respect to cache/blocking strategies and pyramid search primitives (akin to Ansor [zheng2020ansor]).
• Exposing algorithmic alternatives (e.g., Winograd for small filters) as first-class schedule primitives will offer additional avenues for closing the remaining gap to mature libraries.

For Practitioners and Benchmark Designers

• Performance validation must go beyond correctness; correctness-based gates alone admit non-viable kernels with multi-order-of-magnitude slowdowns even in idiomatic code.
• Adoption of the proposed comparability and roofline-based heuristics as screening mechanisms can substantially mitigate fielded regressions and guide optimization campaigns.
• Single-shape, single-GPU, or single-data-type benchmarks are insufficient to guarantee performance portability or universality—kernel evaluation should be recontextualized as an ongoing, iterative process, particularly in LLM-generated code scenarios.

For Automated Code Generation

• The analysis validates and extends recent findings that LLM-generated kernels, even those passing correctness benchmarks, require robust efficiency oracles such as roofline-based metrics to avoid proliferation of low-quality outputs [kernelbench, li2025tritonbench].
• The patterns distilled herein can guide agentic repair loops and automated kernel optimization workflows.

Conclusion

The empirical evidence confirms that the existing evaluation paradigm for GPU DSL kernels, focused solely on correctness, is fundamentally inadequate for quality assurance in performance-critical applications. Severe efficiency collapses are primarily the result of authoring patterns and, to a lesser extent, limitations in code-generation or autotuner coverage. However, with a principled, lightweight screening methodology—anchored in both baseline comparability and hardware-centric roofline analysis—developers can systematically flag and correct most pathological inefficiencies. Future work should focus on improving auto-scheduling coverage, DSL code-generation sophistication, and algorithm-level adaptability to approach the robustness and scalability of vendor-tuned libraries across evolving hardware platforms.

---

References

• "Correct but Slow: An Empirical Study of the GPU Kernel Evaluation Gap in Modern Domain-Specific Languages" (2607.04454).
• [williams2009roofline] Williams, S. et al., "Roofline: an insightful visual performance model for multicore architectures," Commun. ACM, 2009.
• "TritonBench: Benchmarking LLM Capabilities for Generating Triton Operators" [li2025tritonbench].
• "KernelBench: Can LLMs Write Efficient GPU Kernels?" [kernelbench].
• "Ansor: Generating High-Performance Tensor Programs for Deep Learning" [zheng2020ansor].