KernelBench: Can LLMs Write Efficient GPU Kernels?
Authors: Anne Ouyang
Simon Guo
Simran Arora
Alex L Zhang
William Hu
Christopher Re
Azalia Mirhoseini
Venue: ICML 2025 poster (2025)
Keywords: Benchmark
GPU Kernel Design
Code Generation
Primary area: general_machine_learning->evaluation
TL;DR: A benchmark evaluating language models’ ability to generate correct and fast GPU kernels

Abstract: Efficient GPU kernels are crucial for building performant machine learning architectures, but writing them is a time-consuming challenge that requires significant expertise; therefore, we explore using language models (LMs) to automate kernel generation. We introduce **KernelBench**, an open-source framework for evaluating LMs' ability to write fast and correct kernels on a suite of 250 carefully selected PyTorch ML workloads. KernelBench represents a real-world engineering environment and making progress on the introduced benchmark directly translates to faster practical kernels. We introduce a new evaluation metric $\text{fast}_p$, which measures the percentage of generated kernels that are functionally correct and offer a speedup greater than an adjustable threshold $p$ over baseline. Our experiments across various state-of-the-art models and test-time methods show that frontier reasoning models perform the best out of the box but still fall short overall, matching the PyTorch baseline in less than 20\% of the cases. While we show that results can improve by leveraging execution and profiling feedback during iterative refinement, KernelBench remains a challenging benchmark, with its difficulty increasing as we raise speedup threshold $p$.

KernelBench: Can LLMs Write Efficient GPU Kernels?
Anne Ouyang"
Simon Guo
Simran Arora
Alex L Zhang? Azalia Mirhoseini
William Hu Christopher Ré
Abstract
2023; Dao & Gu, 2024a), but their available implementa- tions routinely underperform their peak potential. We are seeing a proliferation of AI hardware (NVIDIA, 2017b; 2020; 2022; Jouppi et al., 2023; Groq; Cerebras; Graph- core), each with different specs and instruction sets, and porting algorithms across platforms is a pain point. A key example is the FlashAttention kernel (Dao et al., 2022), which is crucial for running modern Transformer models - the initial kernel released in 2022, five years after the Transformer was proposed; it took two more years from the release of NVIDIA Hopper GPUs transfer the algorithm to the new hardware platform. We explore the question: Can language models help write correct and optimized kernels?
Efficient GPU kernels are crucial for building per- formant machine learning architectures, but writ- ing them is time-consuee tine-consuming challenge that re- quires significant expertise; therefore, we explore using language models (LMs) to automate kernel generation. We introduce KernelBench, an open- source framework for evaluating LMs' ability to write fast and correct kernels on a suite of 250 carefully selected PyTorch ML workloads. Ker- nelBench represents real-works engincering en- vironment and making progress on the introduced benchmark directly translates to faster practical kernels. We introduce a new evaluation metric fastp. which measures the percentage of gener- ated kernels that are functionally correct and offer speedup greater than an adjustable threshold p over baseline. Our experiments across vari- ous state-of-the-art models and test-time meth- ods show that frontier reasoning models perform the best out of the box but still fall short overall, matching the PyTorch baseline in less than 20% of the cases. While we show that results can improve by leveraging execution and profiling feedback during iterative refinement, KernelBench remains challenging benchmark, with its difficulty in- creasing as we raise speedup threshold
Al engineers use a rich set of information when develop- ing kernels and it is not clear whether language models (I.Ms) can mimic the workflow. They use compiler feed- back, profiling metrics, hardware-specific spees and instruc- tion sets, and knowledge of hardware--fficiency techniques (e.g., tiling, fusion, recompute). They can programming tools ranging from assembly (e.g.. PTX as in DeepSeek-Al (2025)) to higher-level libraries (ThunderKittens (Spector et al., 2024), Triton (Tillet et al., 2019)). Compared to ex- isting LM code generation workloads (Yang et al., 2024a), kernel writing requires massive amount and diversity of information.
We first design environment that reflects the typical AI engineer's 's workflow and supports providing LMs with this rich information. The environment should:
Automate the Al engineer's workflow. The model should have full flexibility to decide which operators to optimize and how to optimize them.
1. Introduction
Al relies on efficient GPU kernels to achieve high perfor- mance and cost and energy savings; however, developing kernels remains challenging. There has been a Cambrian explosion of ML architectures (Tay et al., 2022; Peng et al.,
Support a diverse set of Al algorithms, programming languages, and hardware platforms.
Make it easy to evaluate both performance and functional correctness of LM generations, ideally in a programmatic way. should also capture profiling and exccution infor- mation from generated kernels.
'Equal contribution 'Department of Computer Science, Stan- ford University, Stanford, California, USA DDepartment of Com- puter Science, Princeton University, Princeton, New Jersey, USA. Correspondence to: Anne Ouyang <aco stanford.eduu,, Simon Guo <simonguo stanford.edu>..
We introduce KernelBench to generate and evaluate kernels, which addresses the above considerations. KernelBench tests LM optimizations three levels of Al workloads, each posing different set of challenges:
Proceedings of the 42nd International Conference Machine Learning. Vancouver, Canada. PMLR 267, 2025. Copyright 2025 by the author(s).

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Task Instruction
Evaluate Correctness
Check rucorrect times
Given following model architecture, replace Py Torch operators get speedup by generating inline embodded custom CUDA...
Language Model
Model. forward
Check
Randomized
match?
Inputs
ModelNew fonward
X
clees Model(nn.Moddun
Kemee
class ModelNew(nn. Module):
Measure Performance Benchmark over trtrial times
forward(self,.g x):
CUDA
forward(self, x): operators + custom cuda kernel
calls torch operators
Driver Code
cuda
eager mode torch.compile
Modailiew.
Model. forward
Model
forward
forward
Reference Torch Module
Generated Torch + Compile Inline CUDA
Figure I. KernelBench evaluates LMs' ability to generate performant GPU Kernels. Overview of tasks in KernelBench: KernelBench tasks LMs with generating optimized CUDA kernels for a given target Py Torch model architecture and conducts automated evaluation.
1. Individual operations: We include various AI opera- tors, including matrix multiplies, convolutions, activa- tions, norms, and losses. While PyTorch already uses expert-optimized closed-source kernels, making this a potentially challenging baseline, valuable if LMs can generate open-source kernels for the operations.
for models: while models able to fix execution fail- ures through either reasoning or multiple attempts, they struggle to produce functionally correct code. Further- more, we observe a trade-off between LMs attempting more complex optimizations niche hardware instruc- tions (e-g.. tensor core wmma) and producing error-free kernels. We hypothesize this is due CUDA being low-resource language in open-source training data, only 0.073% of popular code corpus The Stack v1.2 (Li et al., 2023; Kocetkov et al., 2022).
Sequence of operations: We provide problems that con- tain 3-6 individual operations together (e.g. mainloop operator like matmul followed by pointwise operators like ReLU and Bias). This enables evaluating the models' ability to fuse multiple operators.
2. Models demonstrate potential to produce performant kernels via optimizations: We observe a few instances where LMs make algorithmic improvements - e.g.. ex- ploiting sparsity, operator fusion, and utilizing hardware features. We notice more of such instances when we explicitly condition the LM on hardware information (e.g., bandwidth and TFLOP specs) and demonstrations of hardware optimization techniques (e.g., tiling, fusion). While these capabilities remain nascent, LMs do demon- strate potential for generating performant kernels.
3. End-to-end architectures: We select architectures from popular Al repositories on Github including and
pytorch, tuggingface/transs buggiingface/pytooo wwtpee//toooochhhneee wwwpechh-iiggemmmdeee These architectures contain many operations.
Mimicking an Al researcher's workflow, the LM takes Py- Torch reference code as input and outputs an optimized version of the code. Similar the human kernel develop- ment process, our environment enables the LM to iterate with compiler and profiler feedback to refine performance. The LM is free to use any programming language and decide both which parts of the PyTorch code to optimize, and how to optimize them. Our pipeline allows us to feed diverse information to the LMs, including hardware-specift. infor- mation, example kernels, and compiler/profiler feedback.
3. Leveraging, feedback is important for reducing execution errors and discovering faster solutions: By providing execution results and profiler feedback to the LM in context, the kernel quality significantly improves after multiple refinements from 12%, 36%, and 12% in fast, to 43%, 72%, and 18% respectively.
We observe that frontier and open-source models perform poorly out-of-the-box on KernelBench, with OpenAI-ol and DeepSeek-RI R1 matching the PyTorch Eager baseline on < 20% of the tasks. These model-generated kernels greatly suffer from execution errors, functional correctness issues, and are unable to perform cnatform-specine optimizations.
Our findings highlight the technical challenges we need to solve order to adopt LMs for kernel writing. These include but are not limited to: how improve LM perfor- mance in low-resource data regime, and how select from the rich set of information address these challenges, we contribute (1) an open-source framework to study LM kernel generation with a compre- hensive suite evaluation problems and (2) analysis of where current LMs stand and how to realize a future of efficient kernels generated by models.
can provide to models. To
To identify areas for improvement, we conduct series of experiments and analysis, and find that:
Writing functionaly correct kernels remains challenging
2

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
2. Related Works
3. KernelBench: A Framework for Al Kernel Generation
Kernel libraries and compilers. We evaluate existing ap- proaches for kernel programming along the dimensions of automation, breadth, and performance. Mainstream kernel programming libraries like cuDNN (NVIDIA, 2014), CUT- LASS (NVIDIA, 2017a), and Apple MLX (Apple, 2020) are hardware-spee ware-speeiiiee and demand substantial engineering effort from human experts. Other libraries, like ThunderKit- tens (Spector et al., 2024) and Triton (Tillet et al.. 2019), successfully help Al researchers write a breadth of fast and correct kernels (Arora al., 2024; Yang & Zhang, 2024), but still require human programming effort. Compiler-basee tools, like torch.compile (Paszke et al., 2019) and FlexAt- tention (Team Py Torch et al., 2024), automatically provide a narrow slice of optimizations. These methods have pow- erful guarantees in terms of generating provably-correc and robust kernels but are based on fixed operator-fuse fusion or graph-transformate policies (Zheng et al., 2022; Shi et al., 2023; Zhu al., 2022). Previous work on ML--based kernel generation like TVM (Chen et al., 2018) incorporate ML algorithms within single, often sand boxed, component such as schedule generation. In contrast to these efforts, we ask if LMs can automatically generate performant kernels for a breadth of Al workloads.
KernelBench is new framework for evaluating the ability of language models to generate performant kernels for a breadth of Al workloads. In this section, we describe the task format, contents, and evaluation metric.
3.1. KernelBench Task Format
KernelBench contains 250 tasks representing a range of Al workloads, and easily extensible to new workloads. The end-to-end specification for task is illustrated in Figure I and described below.
Task input:
Given an Al workload, the input to the task is a reference implementation written in PyTorch. Mimicking an Al researcher's workflow, the PyTorch code contains a class named Model derived from torch, nn. ww..nooooooooo Module (), where the standard -.. and forward() () functions (and any helper functions) are populated with the AI work- load's s PyTorch operations. We explore alternative forms of imput specifications in Appendix M.
Al algorithms generally operate on large tensors of data. The optimal kernel for a workload depends on the size and data type (e.g., BF16, FP8) of the tensor. Therefore, each task additionally contains functions get () and get get_inite t miiiuutsss (), which specify the exact input tensors that the kernel needs to handle.
LLMs for performance-optimnzze code generation. In the past year, there have been several efforts to build LMs that can automate algorithmic coding (Chen et al., 2021; Shi et al., 2024; Li et al., 2022), resolving GitHub issues (Yang et al., 2024a;b), and domain-specific coding (Yin et al., 2022; Lai et al., 2022). Other works have explored using LMs (Chen et al., 2023; Xia & Zhang, 2024) for program repair and debugging, through models leveraging feedback. While these works focus on producing correct and func- tional code, subsequent works have explored LMs' ability to produce solutions with better algorithmic and asymptotic efficiency (Nichols et al., 2024; Waghjale et al., 2024).
Task output: Given the input, the LM needs to output a new class named ModelNew derived from torch. nn. Module (), which contains custom optimiza- tions. For example, the LM can incorporate in-line kernel calls during the forward () function using the CUDA-C extension in PyTorch.
In order to succeed, the LM needs to identify (1) which operations in the Model class would benefit most from optimizations, and (2) how to optimize those operations. The LM can use any hardware-efficienncy techniques such as fusion and tiling or specialized instructions (e.g.. tensor cores) and any programming library (e.g., PTX, CUDA, CUTLASS, Triton, ThunderKittens) We focus on CUDA in this paper, and explore programming libraries like Triton in Appendix O.
KernelBench focuses on wall-clock efficiency. For LMs to generate perforerfor computing (HPC) code, which requires an understanding of the underlying hardware fea- tures and device instruction set, and common performance characteristics of parallel processors. Existing works in the space of HPC code generation have evaluated LM performance on translating arbitrary code samples from C++ to CUDA (TehraniJamsaz et al., 2024; Wen et al., 2022) or generating nell-knoww, low-level kernels such as GEMMs LValero-Lara et al., 2023; Wijk et al., 2024). Ker- nelBench instead curates set of 250 diverse kernels from real-world, modern deep learning workloads, many of which do not have existing human-writted implementations in other words, solving KernelBench tasks are immediately beneficial for real deep learning workloads. See Appendix for a comparison with popular coding benchmarks.
3.2. Task Selection
The 250 tasks in KernelBench are partitioned into three lev- els, based on the number of primitive operations, or PyTorch library functions, they contain:
Level 1 (100 tasks): Single primitive operation. This level includes the foundational building blocks of AI (c.g.
convolutions, matrix-vector and matrix-matris multiplica-
3

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
tions, losses, activations, and layer normalizations).
2. Performance We compare the wall-clock execution time of Model against Mode NNw using repeated trials to account for timing variations.
Since PyTorch makes calls several well-optimized and often closed-soure kernels under-the-hood, it can be chal- lenging for LMs to outperform the baseline for these prim- itive operations. However, if an LM succeeds, the open- source kernels could be an impactful alternative to the closed-source (e.g.. CuBLAS (NVIDIA, 2023)) kernels.
Comparing LMs on KernelBench Some LMs may gen- erate a small number of correct kernels that are very fast, while other LMs generate a large number of correct ker- nels that are quite slow. Here, we explain our proposed unified metric for ranking LM quality on KernelBench. We elaborate on our process of score design in Appendix I.
Level 2 (100 tasks): Operator sequences. This level includes Al workloads containing multiple primitive op- erations, which can be fused into a single kernel for im- proved performance (e.g., a combination of convolution, ReLU, and bias).
To capture both axes of correctness and performance, we introduce a new metric called fast,. which is defined as the fraction of tasks that are both correct and have a speedup (computed as the ratio of PyTorch wall-clock time to gener- ated kernel time) greater than threshold p. Formally:
Since compiler-base based tools such as the PyTorch compiler are effective at fusion, it be challenging for LMS to outperform them. However, LMs may propose more complex algorithms compared to compiler rules.
N
Σ (correct,, A {speedup, > p}), fast. - N
1
Level 3 (50 tasks): Full ML architectures. This level includes architectures that power popular AI models, such as AlexNet and MiniGPT, collected from popular PyTorch repositories on GitHub.
where fasta is equivalent to the LM's correctness rate, as it measures the fraction of tasks for which the LM code is functionally correct regardless of its speed.
Given the scale of modern models, it is critical to use ker- nels when running training and inference. Unfortunately, it has been difficult for the AI community to generate performant kernels. For instance, it took 5 years from the release of the Transformer architecture (Vaswani et al., 2017) to obtain performant kernels (Dao et al., 2022), let alone today's 's many new architectures. Peak performance kernels for these architectures require algorithmic modifi- cations that are often beyond the scope of a compiler.
By adjusting the threshold parameter p, we enable evalua- tion of kernel performance at different speedup thresholds and capture the speedup distributions. For our evaluations, we focus on p 1 as starting point, with the possibil- ity of increasing p future methods for kernel generation improve. Additionally, using p < 1 for training is valu- able, since PyTorch relies on complex optimized kernels, and matching even a fraction of their performance is still considered beneficial.
We reiterate that each task contains a meaningful set of AI primitive operations architectures, such that LM success on the task can directly lead to real world impact. We provide further details on task definitions and breakdowns in Appendix K.
4. KernelBench Baseline Evaluation
fast, over:
PyTorch Eager
torch.compiin
KernelBench Level
2
3
1
3.3. Metric Design
GPT-4o
4%
5%
0%
18%
4%
4%
OpenAl o1
10%
24%
12%
28%
19%
4%
We describe the evaluation approach for KernelBench and how we compare the success of different LMs.
DeepSeek V3
6%
4%
8%
20%
2%
2%
DeepSeek R1
12%
36%
2%
48%
37% %
2%
Claude 3.5 Sonnet
10%
7%
2%
29%
2%
2%
Evaluation approach KernelBench is an evaluation-only benchmark. We do not provide ground truth kernels for the tasks since we imagine users benchmarking on a variety of hardware platforms (including new platforms), input types, and workloads. However, by design, KernelBench is auto- matically verifiable. Given a task, we randomly generate input tensors of the prescribed shape and precision and col- lect the PyTorch Model output. We can evaluate whether LM generations are correct and fast as follows:
Llama 3.1 3.1 Inst.
3%
0%
0%
11%
0%
0%
Llama 3.1-B05B Inst.
3%
0%
2%
16%
0%
0%
Table I. KernelBench is a challenging benchmark for current LMs. Here we present fast,, i.c. the percentage of problems where the model-generated kernel is faster than the Py PyTorch Eager and t torch. compile baseline (default config) on NVIDIA L40S. The torch. comp le baseline runtime is sometimes slower than Torch Eager - this due to reproducible runtime overhead (nor compile time) that could be significant for small kernlls in Level 1. We focus on PyTorch Eager for the rest of our analysis, but we elaborate other baselines in Appendix B.
1. Correctness We compare the Model output to the LM- generated Mode NNw output. We evaluate on 5 random inputs per problem (detailed in Appendix B).
4

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
4o
OpenAi ol
Deepseek
Deepseek R1
Claude 3.5 Sonnet
Llama 3.1-70b
Liama 3.1-405b
Level 1
Level 2
Level 3
0.6
fast, score
0.2
0.0
0.5
1.5
2
3
0
0.5
1.5
3
0
0.5
1.5
Speedup Threshold (p)
Speedup Threshold (p)
Speedup Threshold (p)
Figure 2. Most LM-generated kernels slow. This figure shows the distribution of the fast, metric the speedup threshold y increases. fasts represents the number of correct kernels regardless of speed, and fast, those achieving at least >> 1x speedup over PyTorch. Increasing p increases difficulty.
In this section, we investigate how a range of LMs perform when evaluated off-the- on Kerneleench and explore their capabilities and failure modes.
Execution Fallure
Functional Correctness
NVCC Error
CUDA Memory
Output Shape Mismaach
Python Error
Runtime Error
Output Value Mismatch
gpt-4o
4.1. One-shot Baseline
openai-o1
We evaluate LMs using a prompt that contains one example of PyTorch Model input and ModelNew output, high- lighting the task format. The example is simple, containing only an add operator (See Appendix C.I). Given this in- context example and the Py task Model to optimize, the LM generates Modell 1New via greedy decoding. We profile the generated code on an NVIDIA L40S GPU, and measure the fast, metric across all problems. Table shows that the L--generated kernels achieves speedup Py- Torch Eager in fewer than 20% of tasks average.
deepseek.Va
deepseck.r1
3aude-3.5-sonnet
llama-3.1-70b
llama-3.1-405b
Percentage of Errors All 250 KernelBench Problems
Figure 3. We categorize failure modes of kernel code into ex- ecution failure and functional correctness. For one-shot baseline, reasoning models generate fewer kernels with execution failures, but models struggle with functional correctness.
4.2. Correctness: Error Analysis
kernels outperform PyTorch across KernelBench levels. Reasoning- based LMs generally outperform the other LMs in providing speedups. We dive into causes of performance degradation in Appendix N.
In Figure 3, we analyze the failure modes of LMs across problems. It can be seen that a large proportion of model- generated kernels are incorrect. To better understand where model-generated kernels fail, we break down their correct- ness issues into execution failures (CUDA/avce) Python compile-time errors, CUDA memory violations, and run- time errors) and correctness errors (output tensor shape and value mismatches). We observe that the reasoning LMs (ol, RI) produce fewer incorrect solutions (< 55%) than other models (>> 70%) due to fewer execution failures. All LMs struggle with functional correctness to similar degree.
4.4. Performance Variations across Hardware
Our one-shot baseline makes assumptions about the derlying hardware, natural question is how analysis of the LM- -Mgeeeertted kernels generalizes across various GPU types. Table 14 and Figure 8 show that kernels outperform- ing PyTorch Eager on NVIDIA L40S in Level achieve sim- ilar speedups versus the baselines on other GPUs. However, on problems in Level LMs exhibit larger variations in speedups across GPUs (Figure 9): DeepSeck R -l-genertted kernels achieve fast, of 36% on NVIDIA L40S but 47% NVIDIA 100G for Level 2. This suggests that one-shot LM-generated kernels may not generalize well across hard- ware. To generate more sprecf-cpecfiic kernels, we explore further in Section 5.2 whether providing hardware-speciine details in-context can improve LM performance.
4.3. Performance: Speedup Distribution
A key point of interest is whether the functionally correct LM-generated kernels outperform the PyTorch baseline. Figure 2 shows the distribution of fast, as p varies, in- dicating the percentage of kernels that are y-times faster than the Py'Torch Eager baseline (the top right of the plot better). At p 1, fewer than 1 15% of L.--enerated
5

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Our analysis reveals that the best models available today
Repeated sampling helps LMs discover more fast and correct solutions. Figure 4 shows that repeated sampling with high temperature improves fast, as k increases across all three levels with both DeepSeek-Ve and Llama 3.1 70B. Notably, on Level 2. DeepSeek- V3 reaches a fast; of 37% with & = 100 samples, compared to just 4% in the one-shot baseline.
struggle generate correct kernels that outperform the base- line PyTorch speeds. LM-generated kernels frequently fail due to simple compiler and run-time errors. Furthermore, it is difficult for LMs to write kernels that perform well across hardware platforms given simple instructions.
5. Analysis of Model Capabilities
Examining the samples, we find that high-temperature sampling helps explore the solution space, increasing the chances of generating error-free kernels with better opti- mizations. However. if a model has a very low inherent probability of solving task, simply increasing the sampling budget has limited impact. For example, DeepSeek- V3 was never able to generate any correct solution for a group of 34 convolution variants in Level 1, even when attempting with 100 samples.
In the last section, we found that KernelBench is a chal- lenging benchmark for today's models. In this section, we conduct case studies to explore opportunities for improve- ment in future models and AI systems.
5.1. Case Study: Leveraging the KernelBench Environment Feedback at Test-Time
As observed in Section 4.2, execution failures are the most frequent failure mode in LM-generated kernlls. The envi- ronment provided by KernelBench allows us to collect rich signals, including compiler errors, correctness checks, and runtime profiling metrics, all of which can be fed back in to the LM to help it resolve kernel failures. To explore how well LMs can use this feedback, we evaluate and compare two baselines: (1) generating multiple parallel samples from the LM per KernelBench task and (2) sequentially gener- ating kernels per KernelBench task by allowing the LM to iteratively refine using the execution feedback.
5, 1.2. ITERATIVE REFINEMENT OF GENERATIONS
The KemelBench environment is well-suited for collecting compiler feedback, execution errors, and timing analysis using tools like the PyTorch profiler as ground-truth signals. We investigate whether leveraging this feedback can help LMs to iteratively refine their generations.
Problem
Model
Generated
Compiter?
Correct?
Performant?
NVCC Compite Issue
Output Mismatch
Execution PyTorch Profiler
5.1.1. 1. REPEATED SAMPLING
Feedbbakk
The KernelBench enviromment enables programmatic verifi- cation of AM-generated kernels, allowing us to collect and evaluate multiple LM generations per task (Brown et al., 2024; Li et al., 2022: Grubisic et al., 2024). We evaluate this repeated sampling approach using fast, @k. which measures the percentage of tasks where the model generated at least one functionally correct kernel that p times faster than PyTorch Eager when drawing & samples.
Figure 5. The KernelBench framework enables models to re- ceive and leverage feedback during iterative refinement. These ground-true signals include NVCC compiler crror messages, exc- cution statistics (correctness and wall clock time), and the PyTorch profiler (operator timing breakdown).
We provide feedhack to the model after each generation in a multi-turn process: after the initial generation, we provide the model with its previous generation G, as well as com- piler/execution feedback F and/or profiler output P over its current generation. We define each generation and subse- quent feedback as a turn, and run this Iterative Refinement process over N turns. For each turn, we measure fast, @N which is the percentage of tasks where the model generated at least one functionally correct kernel that is p times faster than PyTorch Eager by turn N.
Deepseek V3
Llama 3.1 70B
1.0
1.0
Level 1
Level
0.8
Level 2
0.8
Level
Level 3
Level 3
Fast,@k
0.6
Fast,
0.6
0.4
0.4
0.0
0.0
Leveraging execution feedback helps reduce errors and improves overall speedups over time. We examine the fast; behavior at turn N = 10 in Table 2 and find that itera- tive refinement consistently improves performance across models and levels of KernelBench. DeepSeek-RI on Level results in the most notable improvement, where the com- bination of execution feedback and profiler feedback P
10
10a
10
10c
10a
10a
Number of samples,
Number of samples,
Figure 4, Repeated sampling helps discover more correct and performant kernels. As the number of parallel samples in- creases (up to 100), fast, @k improves for both DeepSeek-V V3 and Llama 3.B Instruct across all KernelBench levels.
6

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Level
Level 2
Level 3
Method
Llama
DeepSeek DeepSeek
Llama
DeepSeek DeepSeek
Llama
DeepSeek DeepSeek
3.1 70B
V3
R1
3.1 70B
V3
R1
3.1 70B
V3
RI
Single Attempt (Bascline)
3%
6%
12%
0%
4%
36%
0%
8%
2%
Repeated Sampling (@ 10)
5%
11%
N/A
3%
14%
N/A
1%
14%
N/A
Iterative Refinement G
9%
9%
18%
0%
7%
44%
0%
14%
4%
Iterative Refinement G+E
5%
13%
41 %
5%
5%
62%
8 %
22%
12%
Iterative Refinement G+E+P
7%
19%
43 %
4%
6%
72%
2%
14%
18%
Table Both repeated sampling and iterative improvement enable models to generate more correct and fast kernels compared to baseline: Here present the percentage of problems where the LM- .M-generated kernel is correct and faster than baseline Torch Eager (Fast; in %6 for the two test-time methods, both with the same sample budget of 10 calls. We further compare performance within iterative refinement achieved when leveraging previous Generation G. Execution Result E, and Timing Profiles P.. Note do not repeatedly sample DeepSeek R as its API endpoint does not provide a temperature parameter.
100
Feedback used in iterative Refinement
5.1.3. COMPARING REPEATED SAMPLING AND ITERATIVE REFINEMENT
Last Generation G only
90
Last Generation Execution Result E
Generation Execution Result Profiler
In Table 2, we compare repeated sampling and iterative re- finement given a fixed budget of 10 inference calls. Both methods provide meaningful improvements over the one- shot baseline, with itcrative refinement being more effective in of the cases. However, ultimately we find that the effectiveness of the test-time methods is inherently depen- dent on the quality of the base model. For instance, with repeated sampling, DeepSeek-V2 V3 consistently outperforms Llama-3.1 70B across all three levels. Similarly, with itera- tive refinement, DeepSeek-RI consistently improves using feedback and P. while DeepSeek- -3 and Llama-3.1 70B does not always benefit from having such information.
80
Fast; Score (%)
70
60
50
40
30
20
10
0
2
10
Iteration Turn N
5.2. Case Study: Generating HardwarreEffficiee Kernels via Hardware Knowledge
Figure 6. Iterative refinement with execution feedback E and profiling information enable models to improve kernel gen- erations over turns. as shown in the fast, @NV trajectory of problems where
It is clear that LMs demonstrate limited success at gener- ating hardware-efficient kernels. This is likely due to the scarcity of kernel code in the training data and the fact that the optimal kernel may need change depending on the hardware platform-specifie properties, discussed in Section 4.4. In this case study, we explore providing 1) in-context examples of best- best-practices for kernel engineering and 2) in-context hardware specification details.
DeepSeek-RI on Level The percentage the best generated kernel up to turn N is correct and faster than Py PyTorhh Eager consistently increases as we increase number of turns.
5.2.1. HARDWARE-AWARY IN- IN-CONTEXT EXAMPLES
boosts fast, from 36% to 72% (shown in Figure 6).
Well- kernels often use techniques such as fusion, tiling, recompute, and asynchrony maximize performance. We find that most of the one-shot generated kernels evalu- ated in Section 4 often do not use these techniques. Here, we explore whether providing explicit in-context examples that use these techniques can help the LMS improve their performance on KernelBench. Specifically, we include three in-context examples: GeLU (Hendrycks Gimpel, 2023) using operator fusion, matrix multiplication using tiling (Mills, 2024), and minimal Flash-Atte Attention (Dao et al.,
Furthermore, by examining iterative refinement trajectories, we find that models self-correct more effectively with execu- tion feedback E, fixing issues especially related to execution errors. DeepSeek-RI on Level and 2 can generate func- tional kernel 990 of the tasks within 10 turns of refine- ment (Table 9). However, the remaining incorrect kernels almost always fail due to functional incorrectness, likely because correctness feedback is less granular than execution failure messages. We include successful and failed examples of iterative refinement trajectories in Appendix D.4.
7

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
2022; Kim, 2024) kernel that demonstrates shared memory 1/O management.
6. Discussion
In this section, we discuss qualitative examples of LM gen- erations, and discuss opportunities for improvement.
In-context examples degrade the LM's overall fast, score since LMs attempt more aggressive optimization strate- gies, but result in more execution failures. OpenAl ol ss generations are 25% longer on average using the few-shot examples, compared to the generations produced by Sec- tion 4 baseline. However, among the correct solutions, the LMs apply interesting optimizations: we find that on 77% of GEMM variants in KernelBench Level 1, ol applies tiling and improves speed over the one-shot baseline (although remains slower than Py PyTorch Eager due to the lack of tensor core utilization). On Level 2. ol applies aggressive shared memory I/O management on problems, and is able to outperform PyTorch Eager (See Appendix F).
6.1. Deep Dive Into Interesting Kernels
Here, we discuss a few surprising LM-generated kernels that demonstrate significant speedups over the PyTorch baseline. See detailed examples Appendix D.
Operator fusion GPUs have small amounts of fast-access memory and large amounts of slow-access memory. Fusion can help reduce slow-access I/O costs by performing multi- ple operations on data that has been loaded into fast-access memory. We find that LMs optimize the GELU (2.9x) and Softsign (1.3x) operators by fusing their computations into single kernel. LMs generated kernel that fuses mul- tiple foundational operators matrix multiplication with division, summation, and scaling giving 2.6x speedup. Overall, LMs leave many fusion opportunities on the table, we provide additional analysis on kernel fusion behavior in Appendix L.
5.2.2. SPECIFYING HARDWARE INFORMATION
As discussed in Section 4.4, kernel performance varies depending on the hardware platform. For instance, #lashAttention-22 (Dao, 2024) degrades 47% in hardware utilization going from the NVIDIA A100 to H100 GPU. FlashAttention-33 (Shah et al., 2024), an entirely different al- gorithm, was written for the H100. In this study, we explore whether LMs can use (1) hardware specifications such as the GPU type (H100, A A100, etc.). memory sizes, bandwidths, TFLOPS and (2) hardware knowledge (c.g. definitions of threads, warps, -hocks- streaming multiprocessors) in-context to generate improved kernels.
Memory hierarchy Effective kernels explicitly manage how the limited amounts of fast-access memory (e-g., shared and register memory) gets utilized. In the generated kernels, we found kernels that uses GPU shared memory cosine similarity (2.8x) and triplet margin loss (2.0x) to achieve speedups. We did not find successful usages of tensor core instructions, which are crucial for Al performance.
Models rarely generate kernels that are optimized for the underlying hardware, highlighting room for improve- ment for future models. Certain generations of GPUs (e.g. H100) feature variety of new hardware units and instructions from their predecessors. Providing hardware information does not significantly impact the outputs of Llama 3.1 70B or DeepSeck-V3s V3. Interestingly, we find that a subset of OpenAI ol and DeepSeek-R generated kernels use hardware-specific instructions and optimizations. RI at- tempts to generate warp matrix multiply-accummiin (wmma) instructions (Figure 10) for approximately 50% of the Level I matrix multiplication problems, although most fail to com- pile. Among the functionally correct generations, RI and ol produce 1-3 outliers per level that are 2 2x faster than the Section baselines. Overall, we find that modern LMs are better adjusting their approaches when provided with few-shot examples in Section 5.2.1 than when conditioned on hardware information. Even when explicitly guiding R1 to use architecture-sper speciffc instructions (e.g. wmma and hemcpyyasynn through in-context examples, the model struggled to apply these instructions correctly on simple Level matrix multiplication problems. See Appendix G for more details.
A Algorithmic optimizations Kernels can require algorithmic modifications to better utilize the hardware features. We found one interesting generation for the problem of perform- ing multiplication between a dense and diagonal matrix, where the kernel scales each row (or column), rather than loading the zero-entries of the diagonal matrix, yielding a 13x speedup over PyTorrc Eager.
6.2. Opportunities for Future Work
We show that there is significant room for improvement on KernelBench given the currently available models. First, future work can explore the development of advanced fine- tuning and reasoning techniques, including agentic work- flows. Since CUDA is low-resource language, it would be valuable for future work to open-source more high qual- ity data. Second, LMs generate raw CUDA code in our experiments. However, future work can explore whether generating code using alternative programming abstractions (e.g., provided in ThunderKittens, CUTLASS, Triton, and others) can simplify the generation problem, for instance by making it easier for LMs to leverage tensor core instructions. Third, our evaluation has also been limited to GPUs so far and future work can expand to other hardware accelerators.
8

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
6.3. Conclusion
under No. U54EB020405 (Mobilize), NSF under Nos.
CCF2247015 (Hardware-Awaree, CCF1763315 (Beyond Sparsity), CCF1563078 (Volume to Velocity), and 1937301 (RTML); US DEVCOM ARL under Nos. W91 W91INF-23-2 0184 (Long-context) and W91 INF-21-2-025) (Interactive Human-AI Teaming); ONR under Nos. N000142312633
Our contributions are: (1) We present KernelBench, a frame- work that lays the groundwork for LM-driven kemell opti- mization, and (2) We evaluate diverse set of models and approaches, analyzing their strengths and limitations, and providing insights for opportunities to enhance kernel gen- cration using AI models.
(Deep Signal Processing); Stanford HAI under No. 247183; NXP. Xilinx. LETI-CEA, Intel. IBM, Microsoft, NEC, Toshiba, TSMC, ARM, Hitachi, BASF, Accenture, Ericsson, Qualcomm, Analog Devices, Google Cloud, Salesforce, To- tal, the HAI-GCP Cloud Credits for Research program, the Stanford Data Science Initiative (SDSI), and members of the Stanford DAWN project: Meta, Google, and VMWare. The U.S. Government authorized to reproduce and distribute reprints for Governmental purposes notwithstanding any copyright notation thereon. Any opinions, findings, and con- clusions recommendations expressed in this material are those of the authors and do not necessarily reflect the views, policies, or endorsements, either expressed or implied, of NIH, ONR, or the U.S. Government.
Overall, while most benchmarks eventually saturate, Kernel- Bench is designed dynamically evolve as new Al work- loads arise. Our fast, metric can be adapted over time to measure the speedup threshold (p) over increasingly ad- vanced baselines (i.e.. beyond the PyTorch baseline used in our work). Since PyTorch is cross-hardware platform compatible, the PyTorch-based tasks in KernelBench tasks can be evaluated on every new hardware platform release.
Finally, unlike many benchmarks, success on KernelBench directly maps production value and real-world impacts (lowering costs and reducing energy consumption at scale). These properties ensure that KernelBench will remain valu- able in the ever-evolving Al landscape.
References
Impact Statement
Ansel, J., Yang, E., He, H., Gimelshein, N., Jain, A., Voz- nesensky, M., Bao, B., Bell, P., Berard, D., Burovski, E., Chauhan, G., Chourdia, A., Constable, W., Desmaison, A., Vito, Z., Ellison, E., Feng, W., Gong, J., Gschwind, M., Hirsh, B., Huang, S., Kalambarkar, K., Kirsch, L., Lazos, M.. Lezcano, M., Liang, Y., Liang, J., Lu, Y., Luk, C. K., Maher, B., Pan, Y., Puhrsch, C., Reso, M., Saroufim, M., Siraichi, M. Y., Suk, H., Zhang, S., Suo, M., Tillet, P., Zhao, X., Wang, E., Zhou, K., Zou, R., Wang, X., Mathews, A., Wen, W., Chanan, G., Wu, P., and Chintala, S. Pytorch Faster machine learning through dynamic python bytecode transformation and graph com- pilation. In Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 2, ASPLOS '24, pp. 929-947, New York, NY, USA, 2024. Associa- tion for Computing Machinery. ISBN 9798400703850. doi: 10. (0.114573620666536466666 URL https: /oii. o org/10. /111111///325551 356.522006666666666666636
Optimized GPU kernels can lead to significant energy sav- ings in large-scale machine learning workloads, reducing both computational costs and environmental impact. By providing a framework for Al-assisted performance tun- ing, KernelBench contributes to more energy-efficient AI systems, aligning with global efforts to reduce the carbon footprint of computing infrastructure.
KernelBench does not involve human studies collect user data, climinating privacy concerns. It also avoids proprietary or private code, relying solely on publicly available Github repositories.
Acknowledgements
We thank Aaryan Singhal, AJ Root, Allen Nie, Anjiang Wei, Benjamin Spector, Bilal Khan, Bradley Brown, Daniel Y. Fu, Dylan Patel, Fredrik Kjolstad, Genghan Zhang, Hieu Pham, Hugh Leather, John Yang. Jon Saad-Falcon, Jor- dan Juravsky, Marcel Rød, Mark Saroufim, Michael Zhang, Minkai Xu, Ofir Press, Ryan Ehrlich, Sahan Paliskara, Sahil Jain, Shicheng (George) Liu, Suhas Kotha, Tatsunori Hashimoto, Vikram Sharma Mailthody, and Yangjun Ruan for insightful discussions and constructive feedback in shap- ing this work. We are also grateful to PyTorch, Prime Intel- lect, and Modal for supporting this work.
Apple. Apple ml compute framework (mlx), 2020, URL https://develoy: wwwpllppppp..ppe...pppeeeppp www. coom/metll/.
Arora, S., Eyuboglu, S., Zhang, M., Timalsina, A., Alberti, S., Zinsley, D., Zou, J., Rudra, A., and Ré, C. Sim- ple linear attention language models balance the recall- throughput tradeoff. International Conference on Ma- chine Learning, 2024.
We gratefully acknowledge the support of Google Deep- Mind, Google Research, Stanford HAI, and members of the Stanford SEAMS project: IBM and Felicis; NIH
Brown, B., Juravsky, J., Ehrlich, R., Clark, R., Le, Q. V., Ré, C., and Mirhoseini, A. Large language monkeys: Scaling
9

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
inference compute with repeated sampling, 2024, URL
Graphcore. Grapheore IPU architecture. Online. https: //ww..
battp::/arxiv..e /bsp///////////// abs//407 2407 .000772218777
Cerebras. Cerebras wwfer-scale engine wse architecture. On- line.
Groq. Groq architecture. Online. https:: hettps://groq.com
Grubisic, D., Cummins, C., Seeker, V., and Leather, H. Priority sampling of large language models for compil- ers, 2024. URL https: /awxixxxxvvvoo..a wttp:///gg//bb///////oo// 2402. 18734.
Chen, M., Tworek, J., Jun, H., Yuan, Q., de Oliveira Pinto, H. P.. Kaplan, J., Edwards, H., Burda, Y., Joseph, N., Brockman, G., Ray, A., Puri, R., Krueger, G., Petrov, M., Khlaaf, H., Sastry, G., Mishkin, P., Chan, B., Gray, S., Ryder, N., Pavlov, M., Power, A., Kaiser, L., Bavar- ian, M., Winter, C., Tillet, P. Such, P., Cummings, D., Plappert, M., Chantzis, F., Barnes, E., Herbert- Voss, A., Guss, W. H., Nichol, A., Paino, A., Tezak, N., Tang, J., Babuschkin, I., Balaji. S., Jain, S., Saun- ders. W., Hesse, C., Carr, A. N., Leike, J., Achiam, J. Misra, V., Morikawa, E., Radford, A., Knight, M., Brundage, M., Murati, M., Mayer, K., Welinder, P., Me- Grew, B., Amodei, D., McCandlish, S., Sutskever, I., and Zaremba, W. Evaluating large language models trained code, 2021. URL https: /aaxiv.vvv. org/ / 030703334.
Hendrycks, D. and Gimpel, K. Gaussian error linear units (gelus), 2023. URL pttps::/arxivvorge 1606.08410
Jain, N., Han, K., Gu, A., Li, W.-D., Yan, F., Zhang, T. Wang, S., Solar-Lezama, A., Sen, K., and Stoica, I. Livecodebench: Holistic and contamination free eval- uation of large language models for code, 2024. URL https: attaaaaiivvi...... http:/bb//2230099/// 07974.
Jouppi, N. P., Kurian, G., Li, S., Ma, P., Nagarajan, R., Nai, I.., Patil, N., Subramanian, S., Swing, A., Towles, B., Young. C., Zhou, X., Zhou, Z., and Patterson, D. Tpu y4: An optically reconfigurable supercomputer for machine learning with hardware support for embeddings, 2023, wwwrriii.rr...oeoo worg/bb///////oooooo////o/o//bbr 2204. 01433.
Chen, T., Moreau, T., Jiang, Z., Zheng, L., Yan, E., Cowan, M., Shen, H., Wang, L., Hu, Y., Ceze, L., Guestrin, C., and Krishnamurthy, A. Tvm: an automated end-to-end optimizing compiler for deep learning. In Proceedings of the 13th USENIX Conference on Operating Systems De- sign and Implementation, OSDI' 18, pp. 579-594, USA, 2018. USENIX Association. ISBN 9781931971478.
URL
Kim, P. Flashattention minimal.
Online, 2024.
http::://github.comm/t
Kocetkov, D., Li, R., Allal, L. B., Li, J., Mou, C.,, Ferrandis, C. M., Jernite, Y., Mitchell, M., Hughes, S., Wolf, T. Bahdanau, D., von Werra, L., and de Vries, H. The stack: 3 tb of permissively licensed source code, 2022. URL abs
Chen, X., Lin, M., Schärli, N., and Zhou, D. Teaching large language models to self-debug, 2023. URL https: // //aaxivv.. hbtp://b///334// 05128.
Dao, T. FlashAttention--2: Faster attention with better paral- lelism and work partitioning. International Conference on Learning Representations, 2024.
Lai, Y., Li, C., Wang, Y., Zhang, T., Zhong, R., Zettlemoyer, L., tau Yih, S. W., Fried, D., Wang, S., and Yu, T. Ds- 1000: A natural and reliable benchmark for data science code generation, 2022. URL https: http://axxvvvoooo/e abs/2211 11501.
Dao. Т. and Gu, A. Transformers ssms: Generalized models and efficient algorithms through structured state space duality. International Conference on Machine Learning (ICML), 2024a.
Li, R., Allal, L. B., Zi, Y., Muennighoff, N., Kocetkov, D., Mou, C., Marone, M., Akiki, C.,, J., Chim, J., Liu, Q. Zheltonozhski,, E.., Zhuo, T. Y., Wang, T., Dehaene, O., Davaadorj, M., Lamy-Poirice J., Montciro, J., Shliazhko, O., Gontier, N., Meade, N., Zebaze, A., Yee, M.-H., Uma- pathi, L. K., Zhu, J., Lipkin, B., Oblokulov, M., Wang, Z.. Murthy, R., Stillerman. J., Patel. S., Abulkhanov, D., Zocca, M., Dey, M., Zhang, Z., Fahmy, N., Bhat- tacharyya, U., Yu, W., Singh, S., Luccioni, S., Villegas, P., Kunakov, M., Zhdanov, F., Romero, M., Lee, T., Timor,
Dao. T. and Gu. A. Transformers are SSMs: Generalized models and efficient algorithms through structured state space duality. In International Conference on Machine Learning (ICML), 2024b.
Dao, T., Fu, D. Y., Ermon, S., Rudra, A., and Ré, C. FlashAt- tention: Fast and memory-efficient exact attention with 10-awareness. In Advances in Neural Information Pro- cessing Systems, 2022.
N., Ding, J., Schlesinger, C., Schoelkopf, H., Ebert, J., Dao, T., Mishra, M., Gu, A., Robinson, J., Anderson, C. J., Dolan Dolan-Gavitt, B., Contractor, D.,, Reddy, S., Fried, D., Bahdanau, D., Jernite, Y., Ferrandis, C. M., Hughes,
Dee.Seek-AL. Deepseek-v3 technical report, 2025.
URL https: /github. com/ demn//eeepeeeeeeee aii DeepSeek-V3.
10

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
S., Wolf, T., Guha, A., von Werra, L., and de Vries, H. Starcoder: may the source be with you!, 2023. URL https:/ bthps://aaee /arxiv..vvvv.////i// wwt./r/g//////o////t...//...eee/ 2335 06161.
Shah, J., Bikshandi, G., Zhang, Y., Thakkar, V., Ramani, P., and Dao, T. Flashattention-3: Fast and accurate attention with asynchrony and -oeeprecsi,on, 2024, URL https: / wawxiv.org/abs/2400..1 08608.
Li, Y., Choi, D., Chung, J., Kushman, N., Schrittwieser, J.. Leblond, R., Eccles, T., Keeling, J., Gimeno, F., Dal Lago, A., Hubert, T., Choy, P. de Masson d' Autume, C., Babuschkin, I., Chen, X., Huang, P.-S., Welbl, J., Gowal, S., Cherepanov, A., Molloy, J., Mankowitz, D. J., Sutherland Robson, E., Kohli, P., de Freitas, N., Kavukcuoglu, K., and Vinyals, O. Competition- level code generation with alphacode. Science, 378 (6624):1092-1099) December 2022. ISSN 1095-9203. doi: 10.1 10.11226/siennceaaa11555553 URL http: / org/10. 11..6.11.6666 htt.1/////cieeceeabp////
Shi, Q. Tang, M., Narasimhan, K., and Yao, S. Can lan- guage models solve olympiad programming? 2024. URL https: /tapsi////rvvvvo/// org/ aba///a///// aba//////../ 10952
Shi, Y. Yang, Z., Xue, J., Ma, L., Xia, Y., Miao, Z., Guo, Y., Yang, F., and Zhou, L. Welder: Scheduling deep learning memory access via tile-graph. In /7th USENIX Symposium on Operating Systems Design and Imple- mentation (OSDI 23), Pp. 701-718, Boston, MA, July 2023. USENIX Association. ISBN 978- 978-1-939133--3-2 URL /tt.tttpss///ww.ueennxx wwg.///i.......ooooo.....oooeeee conference/ assdi2//peesentatiion//
Mills, C. J. Cuda mode notes lecture 004. On- line, 2024. post cucd--ooo---otes/eeetteeeeeeeeee
com/
Spector, B., Arora, S., Singhal, A., Fu, D., and Ré, C. Thun- derkittens: Simple, fast, and adorable ai kernels. 2024.
Tay, Y., Dehghani, M., Bahri, D., and Metzler, D. Efficient transformers: A survey. ACM Computing Surveys, 55(6):
Nichols, D., Polasam. P., Menon, H., Marathe, A., Gamblin, T., and Bhatele, A. Performance-aligned lims for gener- ating fast code, 2024. URL https: htttps://arxiv.oo http///axxvvvrr///ee abs/2404 1044..8880 18864.
1-28. 2022.
Team Py He, H., Guessous, D., Liang, Y., and Dong, J. FlexAttention The flexibility of PyTorch with the performance of FlashAttention, 2024. URL https: // pytorch. http:chh..oogoo///o/eeeeooo///// http:///lxxettttttiooo
NVIDIA. cudnn: Gpu-accelerated library for deep neu- ral networks, 2014. URL https: https://develoogg /dee/loeer..eee. nvidia. www./cccuueoooo cudnn.
TehraniJamsaz, A., Bhattacharjee, A., Chen, L., Ahmed, N. K., Yazdanbakhsh, A., and Jannesari, A. Coderosetta: Pushing the boundaries of unsupervised code translation for parallel programming. In The Thirty-eighth Annual Conference on Neural Information Processing Systems, 2024. URL httpt:::////......... /ttps///ppeeeeewwweeeeeee net. forum? id=V6hrg409gg
NVIDIA.
Cuda templates for linear algebra subroutines, 2017a. URL https://github. COM/NVIDIA/ cutlass.
NVIDIA. Nvidia Tesla V100 GPU architecture, 2017b.
NVIDIA. Nvidia A100 tensor core GPU architecture, 2020.
Tillet, P., Kung, H. T., and Cox, D. Triton: an intermediate language and compiler for tiled neural network computa- tions. In Proceedings of the 3rd ACM SIGPLAN Interna- tional Workshop on Machine Learning and Programming Languages, 2019.
NVIDIA. Nvidia H100 tensor core GPU architecture, 2022.
NVIDIA. cuBLAS, 2023. URL ittps://docs.
nvidia, wtm/ cuda/ cubdaa/
Paszke, A., Gross, S., Massa, F., Lerer, A., Bradbury, J., Chanan, G., Killeen, T., Lin, Z., Gimelshein, N., Antiga, L., Desmaison, A., Köpf, A., Yang, F.., Vito, 7., Rai- son, M., Tejani, A., Chilamkurthy, S., Steiner, B., Fang, L., Bai, J., and Chintala, S. Pytorch: An imperative style, terf-performance deep learning library, 2019. URL https:/ / aaa.iiiiii.////vv///aaaa/aaaavva org/ abs / 11199911111 01703.
Turing, A. M. On computable numbers, with an ap- plication to the Entscheidungsproblem. Proceedings of the London Mathematical Society, 2(4 265, 1936. URL http: http:://wwww... /www cs helsinki. fi /u/u/ gionis/ http://00/////// /ttp//////oncoppuuraaeeeeeeeeeee
pdf.
Valero-Lara, P., Huante, A., Lail, M. A., Godoy, W. F. Teranishi, K., Balaprakash, P., and Vetter, S. Comparing llama-2 and gpt-3 1lms for hpc kernels generation, 2023, URL https: / www...vv... org abs 220./ 07103.
Peng, B., Alcaide, E., Anthony, Q., Albalak, A., Arcadinho, S., Cao, H., Cheng, X., Chung, M., Grella, M., Kiran GV, K., He, X., Hou, H., Kazienko, P., Kocon, J., and Kong, J. e. a. Rwkv: Reinventing mnn for the transformer era. Findings of the Association for Computational Linguis- fies: EMNLP 2023, 2023.
Vaswani, A., Shazeer, N., Parmar, N., Uszkoreit, J., Jones, L., Gomez, A. N., Kaiser, L., and Polosukhin, I. Attention is all you need. 31st Conference on Neural Information Processing Systems (NIPS 2017.
2017.

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Waghjale, S., Veerendranath, V., Wang, Z., and Fried, D. ECCO: Can improve model-generated code efficiency without sacrificing functional correctness? In Al-Onaizan, Y., Bansal, M., and Chen, Y.-N. (eds.), Proceedings of the 2024 Conference on Empirical Methods in Nat- ural Language Processing, pp. 15362-15376, Miami, Florida, USA, November 2024. Association for Compu- tational Linguistics. doi: 10.18653/w12204.cmmmpp 859. URL https: ///aaaaaanttoooooooooo wwtp://..rrr2220/0 emnlp emnlp-mas 10pp-pmminn main 859/.
Yin, P., Li, W.-D., Xiao, K., Rao, A., Wen, Y., Shi, K., Howland, J., Bailey, P. Catasta, M., Michalewski, H., Polozov, A., and Sutton, C. Natural language to code generation in interactive data science notebooks, 2022. URL htt.tttt://... arxiv. . org/ //b/////////// aww.bbaa. 2212. 09248.
Zheng, Z.,, Yang, X., Zhao, P., Long, G., Zhu, K., Zhu, F., Zhao, W. Liu, X., Yang, J., Zhai, J., Song, S. L., and Lin, W. Astitch: enabling new multi-dimennsonal optimization space for memory-intensive ml training and inference on modern simt architectures, In Pro- ceedings of the 27th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, ASPLOS '22, pp. 359-373, New York, NY, USA, 2022. Association for Comput- ing Machinery. ISBN 9781450392051. doi: 10.1145/ 3503222.3507723. URL https://doi.orz 3145/3503222.35077
Wen, Y., Guo, Q. Fu, Q. Li, X., Xu, J., Tang, Y., Zhao, Y., Hu, X., Du Z. Li, L., Wang, C., Zhou, X., and auto-parale auto-parallelized program translation. In Chaudhuri, K., Jegelka, S., Song, L.,, Szepesvari, C., Niu, G., and Sabato, S. (eds.), Pro- ceedings of the 39th International Conference chine Learning, volume 162 of Proceedings of Machine Learning Research, pp. 23685-23700. PMLR, 17-23 Jul 2022. URL httpp:: /ttee://rrrrree pwwpeedincgggsssoe mww.m.r.... press/ v1 62 wwn.2b. html.
Chen, Y. BabelTower: Learning
Ma-
Zhu, H. Wu, R., Diao, Y., Ke, S., Li, H., Zhang, C., Xue, J., Ma, L.., Xia, Y., Cui, W., Yang, F., Yang, M., Zhou, L., Cidon, A., and Pekhimenko, G. ROLLER: Fast and efficient tensor compilation for deep learning. 16th USENIX Symposium on Operating Systems De- sign and Implementation (OSDI 22), pp. 233-248, Carls- bad, CA. July 2022. USENIX Association. ISBN 978- 1-939133-28-1. URL //tp:p::/ww/// httttps:/ww..useixx..oom conference/ httped////eeeeeetttio/////////
Wijk, H., Lin, T., Becker, J., Jawhar, S., Parikh, N., Broadley, T., Chan, L., Chen, M., Clymer, J., Dhyani, J., Ericheva, E., Garcia, K., Goodrich, B., Jurkovic, N., Kinniment, M., Lajko, A., Nix, S., Sato, L., Saunders, W., Taran, M., West, B., and Barnes, E. Re-bench: Evaluating fron- tier ai capabilities of language model agents against human experts, 2024. URL https:: /ttps///axxvvvvoee abs/2411- 2411.15114
Xia, C. S. and Zhang, L. Automated Program Repair via Conversation: Fixing 162 out of 337 Bugs for $0.42 Each using ChatGPT. In Proceedings of the 33rd ACM SIGSOFT International Symposium on Software Testing and Analysis, ISSTA '24, pp. 819 831. ACM, September 2024. doi: 10. 10.114553650011233663222 ///xxdoi..ooo 10 11..1....155111111111 360223
URL http:
Yang, J., Jimenez, C. E., Wettig, A., Lieret, K., Yao, S., Narasimhan, K., and Press, O. Swe-agent: Agent- computer interfaces enable automated software engineer- ing, urXiv:2405.15793 2024a.
Yang, J., Jimenez, C. E., Zhang, A. L., Lieret, K., Yang. J.. Wu, X., Press, O., Muennighoff. N., Synnaeve, G., Narasimhan, K. R., Yang, D., Wang, S. I., and Press, O. Swe-bench multimodal: Do ai systems generalize to visual software domains?, 2024b. URL https:// arxiv.orge abs / /.t..221...00 03859.
Yang, S. and Zhang, Y. Fla: A triton-based library for hardware-efficient implementations of lin- ear attention mechanism, January 2024. URL https://github. com/ sustesonglin/ flash-linear-ate
12

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
A. KernelBench Task Example
Here we provide an example task from KernelBench. Each task wrapped class named Model. A task contains two key functions the Model class, init and forward; helper functions are included if necessary. We fix the shape of inputs and vary the numerical values through randomly generated tensors. We provide functions, get_inpute e and get.init.inputs. for generating random parameters for initializing the model and running forward pass, respectively.
PyTorch Reference Architecture:
import torch
import torch nn as nn
class Model (nn, :odule):
=**
Simple model that performs a single matrix miltiplication (C A B) with large
+ dimension .**
def (selt_ (self)
super (Model, self) .off)__iiitt. init_
10
11
def (sele, A: torch. Tenser, B: torch. Tensor) -> torch, Tensor:
12
***
13
Performs matrix multiplication of and B.
14
15
Args:
16
A: Input tensor of shape (K, K)
17
B: Input. tensor of shape (K, N)
18
19
Returns:
20
Output tensor of shape n,, N)
21
***
22
return torch, matmul (A, B)
23
24
M
256
25
256
26
131072
28
def get get_inputs(): ():
29
A = torch. randn (M, K)
30
tarch.randh (K, N) return IA, B1
32
33
def get_init_ingeeeeeeeeee
34
return [] No special initialization inputs needed
We include an example output from model below that attempts to optimize the task specification above. The model not only responsible for generating the kernel code, but is also responsible for generating the peripheral code to integrate the kernel into the PyTorch framework. The evaluation framework will run the forward pass of the model if it is PyTorch operator, common approach is inline the CUDA code.
13

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
import torch
import terch.nn an as mn
from torch.uth uttl filil cpp_exensioon import load_inlime
Define the custom CUDA kernel for matrix multiplication
matmul_ source = cestuul_suurceat
7
#inolude storch/extension.it #inelude <cuda runtime.h>
10
global__ void matmul_kerne (const floate A, const float* B, float C, int M, int N, int
+ K] {
11
int rew blockIdx.y blockDim.y threadIdx.y;
12
int col blockIdx.X blockDim.X threadldx.X/
13
14
it (row < M & col N) {
15
float sum 0.0f;
16
for (int k k ++k) I
17
sum A[row
k] Blk * coll;
18
19
C[row N col] - sum;
20
21
23
: :ensor matmul_cuda : Tensor A, torch: B) {
24
int м A.sizo (0);
25
int N B.size (1);
26
int A.size(1);
28
auto C tarch::zeros ([M, N), A. () ()
29
30
dim3 threadsPerBlock (16, 16);
dim3 numBlocks (N threadsPerBlock.x 1) / tareadsPerBlock.x, (M threadsPerBlockky
1) / threadsperBBlockk)))
32
33
matmul hernnleeeeeel1ccnnlleeeeee threadsPerBlockee (A. da.datt_ttrflloat (), 3.data_ptr<float> ()
C. ptrcflatrr). M, K);
34
35
return C7
36
37
**.
39
ma tmu _ap_____p_urc__e
40
"torch::Tensor hatmul_ccue : A, torch:: ;);"
41
42
Compile the inline CUDA code for matrix miltiplication
44
matmul icad_iniine
45
name="matmul",,
46
48
funetlion==((maatmllLoodddd
49
verbose-True,
50
extra_ccflags-
51
extrra_ldflags=(""]
52
53
54
55
class Newe Module)
56
def _ (self):
57
super ModelNNw, self) init_ 0) self.matmul = matmul
58
59
60
def (srl, torch. Tensor, Tensor) torch. Tensor:
return self. matmul mutaul_ccuaa (A, B)
14

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
B. Evaluation Methodology and Baselines
All evaluations are conducted on a bare-metal NVIDIA L40S GPU with Ada Lovelace architecture unless otherwise stated (such as the device generalization experiments in Section 4.4 and the hardware case study in 5.2). The NVIDIA L40S has 48 GB of HBM memory and operates at 300W. Our environment uses Python 3. 10, PyTorch 2.5. 0+cu12 4, and CUDA 12. 4, which is also where our PyTorch Eager and torch. compile il baselines are derived from.
B.1. Kernel Evaluation Setup
Recall the KemnlBennch task entails a PyTorch reference module Model as baseline, and model model-generated Py PoTcch architecture 1 with custom inline CUDA kernel.
For correctness, we set num correccness to 5, where we check equivalence of output between reference architecture Model and generated architecture with custom kernel ModelNew with 5 randomized inputs. We elaborate on our choice in B.2.
For performance, we measure the wall-clock execution time of nn. monume. forward for both Model and ModelNew We ensure only one kernel being evaluated (no other CUDA process) on current GPU. We warm up for 3 itera- tions and then set num-profile to 100 times which measures the elapsed execution time signaled between CUDA events torch. cuda Event. We take the mean of the 100 trials, and also note its max, min, and standard deviation. While the wall clock time might vary for every trial, we note our coefficient of variation (CV): std/ /mann is consistently < 3%, we use the mean of both measured wall clock time for comparisons.
To compute the speedup of generated architecture over baseline architecture for individual problems, we use the mean for both speedup- = Todeel MoodeNNew- For example, if Model = 2 ms and Node = 1 ms, we have a 2x speedup with the newly generated kernel. We compare this speedup with our speedup threshold parameter p (as explained in section 3.3) to compute fast, scores.
B.2. Correctness Analysis Varying Number of Randomly Generated Inputs
Checking equivalence of programs in a formal sense is undecidable. "The Halting Problem" (Turing, 1936) states that it impossible decide, in general, whether given program will terminate for every possible input. This problem naturally extends checking equivalence because in order to check whether two programs equivalent, it is necessary to check their behavior for all inputs, including cases where one or both programs may not terminate. Since determining whether a program halts on a given input is undecidable (the Halting Problem), checking equivalence also becomes undecidable.
Approximate or heuristic methods are often used in practice for checking program equivalence. Random testing is the most common practical approach, where the program is run with sets of randomly chosen imputs, and their outputs are compared. Random testing is particularly effective for Al kernels, where control flow is simpler and the focus primarily on numerical correctness. By using diverse inputs, can uncover errors in computations or memory handling with high probability.
We use five sets of random inputs for correctness, which is good tradeoff between the ability to catch errors and efficiency. In an experiment with 100 generated kernels, the results were as follows: 50 kernels were correct (all 5/5 and 100/100), 19 had output value mismatches (19 0/5 and 0/100), 4 had output shape mismatches, 10 encountered runtime errors, and 17 had compilation errors. Notably, the 0/5 and 0/100 failures indicate that no partial correctness was observed.
B.3. Distribution of Model Performance for One-Shot Baseline
Here we examine the quality of (functionally correct) kernel generations across a wide variety of models. Figure 7 shows the distribution of speedups for various kernels across different levels and models. The median speedup for both Level and Level are less than 1, and the median speedup for Level 2 is only slightly above one. Level has the most significant outliers, in one case showing a speedup greater than 10. We explored some of these outlier cases in greater detail in Section 6.
Reasoning-optimmieed models (OpenAI-ol and DeepSeek-R1) perform the best of out-of-the-box across all levels. These models demonstrate superior kernel generation capabilities, particularly excelling at Level 2 tasks (which mainly involve kernel fusion). In contrast, Llama 3.1 models (both 405B and 70B) perform poorly regardless of model size, suggesting that larger models do not necessarily guarantee better results for this task. DeepSeck-BB while strong at Level I and 2, suffers significantly at Level 3, often generating incorrect kernels.
15

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Speedups
10
Box and Whisker Plot of Speedups
5
Levels
Level 1
Level
Speedups
3 Speedups
claude-3 3.5 5-sonnetttttt (539)
opt-4o (39%)
deepseek-1 (67%)
openai-o1 1o (35%)
owene 2.5 ----ooda 320 (41%%)
deepseek-v3 (59))
claude-3 3.5 5-sonnet siouue-3.--oonntt (10%)
gpt-4o (9%)
deepseek-11 (62%)
openai-ol o1 (56%)
qwen- 2.5-coderrd3222 32b (10%)
deepseek-vs (30%)
Clauuee3.5-somnet (12%)
gpt-4o (349%)
Ioeepeeeeo 2 (9%)
deepseok-V3
claude-3 Souude-3.5--000 5-sonnet
Models
Figure 7. A box and whisker plot of the speedup relative to Torch Eager of (correct) kernels generated by various models in the one-shot baseline setting. also write percentage of correctly generated kernels next to model name. observe among most models, the median speedup for correctly generated kernels is below 1.
B.4. PyTorch Baselines
PyTorch offers two common execution modes: Eager and torch. compile. Aside from the results shown in Table 1, all performance analysis is evaluated against PyTorch Eager.
PyTorch Eager is the default execution mode of PyTorch, which dynamically executes computation by invoking calls to highly optimized closed-source kernels.
PyTorch Compile or torch. compile uses rule-based heuristics over the underlying computation graph during an initial compilation phase and invokes various backends perform optimizations like kernel fusion and graph transformations. In Table 1, our performance baseline for torch. compile assumes the default configuration using Py Torch Inductor in default mode. Furthermore, we exclude the torch. compile compile time in our timing analysis, as we are only interested in the raw runtime behavior. torch. compile features multiple other backends and configurations, which we deseribe in Table 3.
observe that torch. . compile baseline runtime is generally faster on Level and of KernelBench reference problems compared to PyTorch Eager, mostly due the availability of graph-level optimizations like operator fusion. However, Level 1 problems, torch. compile exhibit higher runtimes than PyTorch Eager, which can be attribute to empiricalll--eproddcibbee runtime overhead for torch. compi le (not compile time) that is significant for small kernels.
Other torch. compile backends. In Table 4, we show more one-shot bascline results for fast, against some of the other torch. me baselines. We note on some other configurations fast, drops especially for Level 2, as the torch. compile ile backends apply more aggressive optimization (at the cost of extra compile- time overhead, which we do not measure). Due the variability of torch. compi e across configurations, we focus our analysis on PyToorch Eager.
16

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Configuration
Backend
Mode
Description
PyTorch (Eager)
Standard PyTorch cager execution
Torch Compile
inductor
default
Default torch. compile le behavior
Torch Compile
inductor
overcee
Optimized for reduced overhead
Torch Compile
inductor
max-autotune
Maximum autotuning enabled
Torch Compile
inductor
maz--attouune-nn-uuuagaaphs
Maximum autotuning without CUDA graphs
Torch Compile
cudagraphs
CUDA graphs with AOT Autograd
Table 3. Configurations and modes for PyTorch exccution and optimization backends.
torch.come corpico default
max-autotune no-cudagraphs
fast, over:
cudagraphs
max-autotune
reduce-overhead
KernelBench Level
2
3
1
2
3
2
3
2
3
Claude 3.5 Sonnet
29%
2%
2%
31% %
7%
2%
31 %
2%
0%
29%
2%
2%
31%
2%
0%
DeepSeek V3
20%
2%
2%
21%
4%
20%
21%
2%
2%
20%
2%
2%
21%
2%
0%
DeepSeek R1
38%
37%
2%
42%
52%
0%
42%
29% %
0%
38 %
32%
4%
42%
28 %
0%
GPT-4o
18%
4%
4%
22%
6%
6%
21%
4%
2%
18%
3%
4%
21%
4%
0%
Llama 3. 3.1-70B Inst.
11%
0%
0%
12%
0%
0%
12%
0%
0%
11%
0%
0%
12%
0%
0%
Llama 3.1-405B Inst.
16%
0%
0%
16%
0%
4%
16%
0%
0%
16%
0%
0%
16%
0%
0%
OpenAI 01
28%
19%
4 %
33%
37%
26 %
34%
8%
%
30%
19%
6%
34%
8%
2%
Table 4. We compare KernelBench torch. compile le baseline runtime across various configurations, all measured on NVIDIA L40S, in addition to what is showed Table
C. Experiment Prompting Details
provide details for the prompting strategies and associated sampling strategies used Section and Section
C.1. One-shot Baseline Prompt
For the one-shot baseline as shown in Section 4.1, we want to examine each model's out-of-the-box ability to generate kernels by providing the minimum set of information while ensuring the instructions and output format are clear. We query each model with the following prompt and pair of in-context add examples (the PyTorch reference add and its CUDA kernel counterpart using inline compilation) to provide the output format. We sample the model with greedy decoding ensure deterministic output, which is setting temperature 0.
You write custom CUDA kernels to replace the pytorch operators in the given architecture to get speedups.
You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others 6 unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.
10
11 Here\'s an example to show you the syntax of inline embedding custon CUDA operators in
12 torch: The example given architecture is:
13 111
14 import torch
15 import torch.m. as nn
16 import torch.nn. functional
17
18
19
class Model (nn. Monule)
20
def _niit____ss (self) -> None:
21
super () ()__)nn__i__e___e_______________ init_iin
17

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
23
def forwar a, b):
return +
25
26
27
def get get_inputs()
randomly generate input tensors based on the model architecture a randn 128) cuda () 29
30
b = torch.randa (1, 128) cuda
31
return [a, b]
34
def
35
randomly generate tensors required for initiali iitiillizatii zation based the model architecture
36
return
37
38
39
The example new arch with custom CUDA kernels looks like this:
40 ...
import torch
42
import torch.nn nn as mn
43 import torch. nn. on. functional as
44
from torch.utils. extension import Load inline
Define the custom CUDA kernel for element-wise addition 46
47
el sementwise_add_soee
48
#include storch/extension.hh
49
#include <cuda_runtime.hs
50
51
kernel global__ void elementwise add kernel (const float= const floatx b, float* out, int
< size) {
52
int idx blockIdx.x blockDim.x threadIdx.x;
if (idx size) {
54
[idx] = [[ixx] b[idx];
55
56
57
torch: :Tnnsor saaneettwss__ad_ccoe (torch: Tensoor a, torch:: iensor
59
auto size = a. a.numel();
60
auto out = torch: izeros (a);
61
62
const
alock_size 256;
const int num_blocks (size block_size - 1) / block_size;
64
65
elementwise add www.adaeeeeeeenencccoco......... float> (),, out dut.data__nee pnttdata__ttrfllat (), size);
block_size>>> (a.
(), b.data_ptar
66
67
return out;
68
69
70
#1 eeemmntwise_adde aad__ppp_sour___ "torch:: :ensor elementwi [ se_add_cud__e (torch: :Tensor a, torch::
4 Tensor b)
72
73
Compile the inline CUDA code for element-wise addition
elementwise add load inline
75
name=' came==/elemmetvvse__ddd!
76
77
78
79
verbose=True,
80
extra_cflags-((1]
81
extra_ldflags=(1
82
18

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
84
class ModelNew (nn (nn.uuuulel :
def _ (self) -> None:
86
super iniit_
87
self sll.feeleenttiseeeee add = elementwise_add
88
89
def (sel,, a, b):
90
return
www.:ddoooootttttsdddeeodooo (a, b)
91
92
You are given the following architecture:
95
<PyTorch reference architecture for specific Kernel Hench Problem
96
97
Optimize the architecture named Model with custom CUDA operators! Name your optimized output architecture ModelNex. Output the new code in codeblocks. Please generate real code, NOT pseudocode, make sure the code compiles and is fully functional. Just output the new model code, no other text, and NO testing code!
99
100
C.2. Repeated Sampling Prompts
For repeated sampling, the same prompt that used for the one-shot baseline in Appendix C.I. We used the same sampling temperature described in (Brown et al., 2024) as they allow sample diversity while ensuring quality. Specifically we use temperature 1.6 for Deepseek- V3 and temperature 0.7 for Llama 3.1-70B.
C.3. Iterative Refinement Prompts
For iterative refinement, we start with the same initial prompt that we used for the one-shot baseline in Appendix C.1. A limitation of our experiments is that we sample with temperature- 0 to focus on the effect of iterating based on feedback rather than introducing variability, On subsequent generations, we prompt the model with the following template depending on the feedback expects:
<Initial prompt from one-shot baseline for specific fic XernelBench problem.>
Here is your latest generation: <Previously generated kernel
Your generated architecture ModelNew and kernel was evaluated on GPU and checked against
+ the reference architecture Model.
Here is your Evaluation Result:
<Raw Compiles 1er and Execution Feedback from stdout> 10
11
<'if correct:'>
12
Your kernel executed successfully and produced the correct output.
13
Here is your wall clock time: (runtime] milliseconds
14
15
<Profiler information if used and cerrect.>
16
17
Name your improved output archi ModelNex. Output the code in codeblocks.
Please generate real code, NOT paeudocode, make sure the code compiles and is fully functional. Just output the new model code, no other text, and NO testing code!
For the compiler and execution feedback, we handle timeouts and deadlocks explicitly with "Your kernel execution timed out", but do not provide any other information.
C.4. Few-Shot in Context Prompts
For Few-Shot experiments as outlined in Section 5.2.1. We provide more details about the in-context example in Appendix F. We sampled these experiments with temperature
<Initial Task prompt from one-shot. basel ine for Instructions
19

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
<initial pair Reference PyTorch and CUDA kernel equivalent for exanple add kernel from 4 one-shot baseline for Instruction>
4 Example <i>
Bere is example architecture <PyTorch reference architecture for No. in-context example>
Here is an optimized verison with custer. CUDA kernels: <PyTorch architecture with Custom CUDA Kernel for No. i in-context example>
10
up to number of in-context sample times
12
13
14
Task:
15
Here is an example architecture:
16
17
<PyTorch reference architecture for specific KernelBench Problem>
18
19
Name your new improved output architecture ModelNew. Output the new code in codeblocks.
Please generate real code, NOT pseudocode, make sure the code compiles and is fully functional. Just output the new model code, no other text, and NO testing code!
C.5. Hardware Case Study Prompts
Here we provide hardware information. This is used in Section 4.4 and elaborated more in G, sampled with temperature 0.
<Initial Task prempt from one-shot baseline for Instruction <Initial pair of Reference PyTorch and CUDA kernel equivalent for example add kernel from 4 one-shot baseline for Instruction>
Here is some information about the underlying hardware that you should keep mind.
The GPU that will run the kernel is NVIDIA <GPU NAME>.
7
have
CDDR6 with ECC of CPU Kemory.
We have <x> CB/s of Memory Bandwidth.
10 We have <x> of RT Core Ferformance TFLCPS.
11 We have <x> of FP32 TELOPS.
12 We have <x> of TF32 Tensor Core TELOPS.
13
14 We have <x> FP16 Tensor
TELOPS.
Tensor Core TELCPS.
15 We have <x> of Peak INTS Tensor TOPS.
16 We have <x> of Peak INT4 Tensor TOPS.
17
have <x> 32-bit registers per SK of Register File size.
18
We have
Maximum nunber of registers per thread.
19 We have <x> of Maximum number of thread blocks per SM.
20 We have <x> KB of Shared memory capacity per SM.
21
We have <x> KB of Maximum shared memory per thread block.
22
24
25
Here are some cencepts about the Ceu architecture that could be helpful:
26
27
Thread: thread single execution unit that can run a single instruction at time.
Thread Block: A thread block is group of threads that can cooperate with each other.
29
Shared Memory: Shared memory is menory space that can be accessed by all threads in
+ thread block.
30
accessed by single thread. Memory Hicrarchy: Memory hicrarchy a pyramid of memory types with different speeds
Register: A register is small memory space that
+ and sizes.
32 Memory Bandwidth: Memory bandwidth is the rate at which data can be road from or stored
+ into memory.
33 Cache: Cache is small memory space that stores frequently accessed data.
HRM: HBM is high-bandwide memory technology that uses 30-stacked DRAM.
20

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
35
36
Here are some best practices for writing CUDA kernels on GPU
37
38 Find ways to parallelize sequential code.
39
Minimize data transfers between the host and the device.
40
Adjust kernel launch configuration to maximize device utilization.
41
Ensure that global memory accesses are coalesced.
42 Minimize redundant accesses to glebal memory whenever possible.
43 Avoid long sequences of diverged execution by threads within the wazp. #Wle added this to reference the specific GPU arenitecture
Use special ized instructions based on the specific fic GPU architecture
46
47
You are given the following architecture:
48
49
<PyTorhh reference architecture for specific KernelBench Problem>
50
51
Name your new improved output architecture ModelNew. Output the new code in codeblocks.
Please generate real code, NOT pscudocode, make sure the code compiles and is fully functional. Just output the new model code, no other text, and NO testing code!
D. Kernels of Interest
In this section we provide examples of interesting or notable kernel generations. We first expand on the discussion in Section 6, where we defined the following categories optimizations: algorithmic optimizations, operator fusion, and using hardware features.
D.1. Algorithmic Optimizations
13x Speedup on Level Problem 11 by Claude-3.5 Sonnet
The original torch operator is torch. diag (A) B, multiplying diagonal matrix formed from the vector A with the matrix B. The model identifies optimization in the special case of diagonal matrix multiplication, where the diagonal matrix doesn't t need to be explicitly construeted. Instead, each element of the vector is directly multiplied with the corresponding row in matrix B, significantly improving performance:
globol vold MHag_matmu _____amm___keellll _ernel
const fooat diag,
const floats mot,
float
const Int
const int M)
const int blockidx.y blockDim.y threadidx. const int co. "blockdxx.x blocknim.x threadráx.w;
10
11
(rew
12
diag
col);
13
14
D.2. Kernel Fusion
2.9x Speedup on Level 1 Problem 87 by DeepSeek-V. V3 GeLU reference in torch:
0.5 x + (1.0 torch, tanh (math.sqrt. (2.0 math.pi) (x + 0. 044415 torch.por (x, 3.0))
Optimized version fuses in a single kernel. There is also small constant folding optimization, instead of computing
21

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
math. sqrt (2.0 math.pi) repeatedly, the kernel uses the precomputed value 0. 7.778845008288554
global void gelu_kernel (const float* X, float= out, int size)
int idx blockIdx.x blockDim.x threadIdx.X;
if (idx < size) {
4
float x_val x [iixx];
float cdf 0.5f (1.0f tanhf (10. 7007998955688226654 (x_val 0.044715E x_ vav x_val))));
out [idx] = x_v * edf;
1.3x Speedup Level 1 Problem 29 by Claude-3.5 Sonnet SoftSign reference in torch:
torch. abs (x)
Fused kernel:
global_ veid softsign_kernel (const float input, float* output, int size) int idx blockIdx.x blockDim.x threadIdx.
if (idx size)
float input [idx];
float abs_x abs (x)
output [idx] = (1.0f + abs ;bs_x);
6x Speedup on Level 2 Problem 13 by Claude-3.5 Sonnet Sequence of operators in torch:
x torch.c matmul (x, self, welf.weight.. Gemm
X / 2 Divide
torch.sum (x, dim-1, keepdim-True Sum
self.scaling factor $ Scaling
Fused kernel:
22

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
glopal__ void fused ous ops__pps___seee_e kernel
const float* input,
const float weight,
float* cutput,
const float scaling_fasten
const int batch_size
const int input_size,
const int hidden_size
{
10
// Each thread handles one element in the batch
11
const int batch_idx = blockIdx.x blockDim.x + threadldx.x)
12
13
if (batch_idx batch_size) {
14
float sum 0.01;
15
16
1/ Compute matmul and divide for this batch element for (int h 0; h < hidden_size; h+4) fleat elem 0.0f;
17
18
19
fint i 0; 1 imput_sizer i++) {
20
elem += input [batch_ids input_siiii size i]
21
weight (h input_size i];
22
23
// Divide by as we go sum += (elem / 2.0f);
24
25
26
27
// scale and store final result output [batch_idx] sum scaling_factor,
29
30
Despite this good example, overall Level generated kernels show insufficient fusion. It is expected that Level problems be expressed in single fused kernel.
1.9x Speedup on Level 3 Problem 49 by OpenAT-ol
We have a hypothetical architecture of the attention mechanism where the softmax is replaced with a ReLU.
att (q @ (-transposcc-2, -1)) + (1.0 / math. sgrt (k.size (-1)))
att att.masked_fil1 fft......a.......aaaal..... :T] = 0, float (* ('-inf() att F.relu(att)
The model found an optimization that fuses the scaling, masked fill, and ReLU but not anything else, resulting in a modest improvement of 1, 1.9x.
23

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
glopal__ void fused_mas masked fill scales kerue_ const fleat+ restrict__ att, const float restrict__ bias, float+ restrict output, int total_elems, float scale, int T, float negative_infinity
{
10
int idx blockIdx.x + blockDim.x threadIdx.X;
11
if (idx < total_slems) {
12
float val att att [idx] scale;
13
int bias_idx idx (T T);
14
it (biasibias idx 0.01)
15
val negative infinity;
16
17
if (val 0.0f)
18
val 0.0f;
19
20
output [idx] = val;
21
22
D.3. Hardware Features
2.8x Speedup on Level 1 Problem 96 by OpenAI-ol Torch reference for Cosine Similarity Loss
cosine_sim terch. nn. functional www.etoooa....cii....eiiiiaeaee. wwwine_similaritty......ooosooee (predicttoon targets, dim=1) return mean (1 cosi cosine_sim,
The generated kernel uses shared memory for reduce redundant global memory accesses, improving data locality and increasing overall performance. This is a moderately complicated kernel with synchronization points and reductions that would be tricky for humans to get right.
global void cosine wirimaaiiiiiii cass_ kernel [css_kerne]]
const float* restrict predictions,
const float* restrict targets,
4 float* restrict_... losses,
const int batch_size, input_size
const
{
// Each black handles one sample in the batch int sample_ide _ix blockIdx.x; if (sample_idx
10
batch_size) return;
12
// Shared memory for reductions extern shared__ float sdatall;
13
14
15
// Pointers to data for this sample
16
const float pred predictions + sample_idx input_size; const float* targ targets + sample_ids input_sizer
17
18
19
// Intermediate float thread_dot 0.0f; float thread_pred_nomm_.s. 0.0f;
for dot preduct and
20
21
22
float thread_targ_norm__s
0.0f;
23
24
for (int idx threadIdx.x; idx input_sizez idx blockDim.x) float p pidx]
24

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
26
float [idg] thread_dot +- t;
27
28
thread__rre__nom_ss +=
29
thread_targ_nomm sq += t;
30
32
// Reducticn for dot product sdata (tatatthreadddk.k thre
33
34
synethreads ;);
for (unsigned int blockDim.x 2; 0;
36
if (threadidx, < s)
37
sdata [threadd +- sdata [thread dx.x + sli
38
39
syncthreads())
41
float dot jot_uurouct sdata(0];
42
43
// Reduction for pred_norm_sq sdata[threadddk.xx throad_prooLnozn_sq
44
syncthreads ;
46
for (unsigned int = blockDim.x 2; > 0; >>= 1)
47
if (threadIdx.x s)
48
sdatalthreadIdx.xx += sdata (treeddtrrd sl;
50
syncthreads(1;
51
}
52
float norm_pred sqrtf (sdata [0] 1e-8f);
53
54
// Reduction for targ norm sq sdata (teteaathhdd thread thread_targ_nont Sarrg_norr__s__ sqj
55
56
syncthreads().
57
for (unsigned int blockDim.x 2; 0; >>=
58
if (threadIdx. s) sdata(threadIdx. x] += sdata (threadIdx.x sl;
59
60
61
syncthreads
62
float perm_targ sqrtf (sdatal0) 1e-8f);
64
65
it (threadIdx.x == 0)
66
float cosine_sim dot_product / (norm_pred norm_targ 1e-8f);
67
losses [sample_idx] 1.0f cosine_sim;
69
1.9x Speedup on Level 1 Problem 98 by Deepseek-Ri Torch reference for Cosine Similarity Loss
self. 1oss_osss_____ in = torch.nn. nn. www..rrppeeMaagginooo. (margin=mazgin) self. loss in (anchor, positive, negative)
Another example of generated kernel using shared memory:
25

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
glopal__ void triplet
triplet_margin www.mmarrnnoooo Loss (
const float+ anchor,
const float positive,
const flcat negative,
float+ icsses,
float margin,
7
int feature_size)
extern shared float shared_surs ;];
10
11
int batch_idx blockIdx.x;
12
int tid threadldx.xi
13
14
int offset batch idx feature_ size;
15
16
const float* anchor + offset;
17
const float* p positive offset;
18
const float negative offset;
19
20
float a_P_sum 0.0f;
21
float a_n_sum 0.0f;
22
23
int stridc blockDim.X;
24
for (int i tid; < size; i += stride)
25
float diff_ap ali] - plil;
26
ap_ssu += diff_ap diff_ap; float diff_an ali] n[il;
27
a_n_sum +- diff diff_an;
29
30
31
shared sums [idd sum; shared sums bblockDim.x tid] a n. a_m__um;
32
33
34
syncthreads ;
35
36
for (int s blockDim.x / 2; 0; >>= 1)
if (tid < s)
38
shared_sums [tid] +- shared_sums [tid + s];
39
shared_sume sums (blockDim.x + tid] += shared sums (blockDim.x tid + sl:
40
(y
43
44
if (tid -- 0)
45
float d ap sqrtf (shared sums(0]);
float d an sqrtf (shared surs [blockDim.x));
47
losses (baechsbachh_icee idx] fmaxf ad_ap - an + margin, 0.0g);
48
49
D.4. Iterative Refinement Examples
D. ITERATIVELY TRYING NEW OPTIMIZATIONS
We provide an example kernel that iteratively improves on its existing generation. the following example, the model attempts optimizations incorrectly, fixes them, and continue to attempt new optimizations, improving its kernel to faster than the torch. compi le baseline (1.34ms) but short of the Torch Eager baseline (0.47ms).
Level 1, Problem 63: 2D convolution with square input and square kernel. DeepSeek-RR with Execution and Profile Feedback
In this example, we see a speedup average kernel runtime from initial generation, where the model repeatedly
26

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Turn
Turn
Turn 3
Turn 4
Turn 5
Turn 6
Turn 7
Turn 8
Turn 9
Turn 10
Compiles? Correct?
X
X
x
X
V
X
v
X
V
Runtime (ms)
9.1
1.57
1.83
1.43
1.13
1.46
Table Iterative refinement trajectory of DeepSeek-RI with execution feedback E and profiler feedback P on Problem 63, Level 1. Torch Eager baseline runs in 0.47ms and torch. compile runs in 1.34ms.
(incorrectly) refines its kernel, fixes the compiler issues using feedback, then continues attempt more optimizations. The first big jump in performance (Turn Turn 3) occurs because the model decides launch thread blocks along output channel dimension, when it originally computed these elements sequentially. The model then attempts to use shared memory in Turn 5, and continues using it, along with texture cache memory with the _.1dq instruction in Turns and 8.
D.4.2. LEVERAGING FEEDBACK TO CORRECT KERNEL CODE
Level 2, Problem 73: 2D Convolution with a BatchNorm and a scale factor. DeepSeek-RI with Execution Feedback We provide an example of kernel that the model struggles to generate correctly, and produces a correct kernel after iterative refinement using execution feedback.
Turn
Turn 2
Turn
Turn
Turn
Turn 6
Turn 7
Turn 8
Turn 9
Turn 10
Compiles?
Correct?
x
X
X
X
X
X
X
X
X
Runtime
3.16
Table Iterative refinement trajectory DeepSeek-RI with execution feedback on Problem 73, Level Torch Eager baseline runs 0.1 0.105ms and torch. comp! T runs in 0.156ms.
In the above example, the model continually produces either the wrong output tensor shape the wrong values and iterates on its kernel using this feedback until the final turn, where it generates functionally correct, albeit non-performant kernel. We provide another example below that explicitly leverages compiler feedback to fix compiler errors:
Level 2, Problem 23: 3D Convolution with GroupNorm and return the mean across all but the batch dimension. DeepSeek-RI with Execution Feedback
Turn 1
Turn 2
Turn 3
Turn 4
Turn 5
Turn 6
Turn 7
Turn 8
Turn 9
Turn 10
Compiles? Correct?
X
X
V
x
/
/
X
x
Runtime
11.4
1.36
1.39
1.33
Table Iterative refinement trajectory of DeepSeck-R1 1 with execution feedback on Problem 23, Level 2. Torch Eager haseline runs in 1. 1.29ms and torch. compils le runs 0.719ms.
In the above example, the model attempts to use the CUB library, but incorrectly invokes function calls. The model is then able to correct these errors and write a slightly faster kernel in Turn 8.
D.4.3. ITERATIVE REFINEMENT NEVER FIXES THE ERROR
Level 1, Problem 54: 3D Convolution square input and square kernel. DeepSeek-R1 with Execution and Profiler Feedback
This problem is particularly interesting because no model is able to consistently produce functional code for this kernel, even with different forms of feedback and profiling information. Interestingly, the example before is an arguably more difficult version of this kernel that fuses the 3D convolution with another operator, and the same model is able to generate functional code for this task. the example above, the model consistently makes the same mistake and continually generates
27

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Turn 1
Turn 2
Turn 3
Turn 4
Turn 5
Turn 6
Turn 7
Turn 8
Turn 9
Turn 10
Compiles? Correct?
V
v
X
X
X
X
X
X
X
Runtime
Table Iterative refinement trajectory of DeepSeek-RI with execution feedback E and profiler feedback / on Problem 54, Level 1. Torch Eager baseline runs in 4.47ms and torch. compile runs in 4.67ms.
a functionally incorrect kernel with the same value errors.
E. Iterative Refinement on Correctness
Here we show that fast, across iterative refinement 5.1.2 configurations at turn budget of N 10 compared to one-shot baseline 4.1. We find that models self-correct more effectively with execution feedback E, fixing issues especially related to execution errors. Notably, DeepSeek-R. on Level and 2 can generate functional kernel on >90% the tasks given 10 turns of iterative refinement. However, the remaining incorrect kernels almost always fail due to functional incorrectness, likely because correctness feedback less granular than execution failure messages
Level
Level 2
Level 3
Method
Llama DeepSeek DeepSeek
Llama
DeepSeek DeepSeek
Llama
DeepSeek DeepSeek
3.1 70B
V3
R1
3.1 70B
V3
RI
3.1 70B
V3
RI
Single Attempt (Baseline)
26%
43%
67%
0%
6%
62%
0%
30%
8%
Iterative Refinement (w G)
27%
48%
72%
2%
7%
67%
36%
14%
Iterative Refinement (w G+E)
40%
53%
95%
7%
8%
85%
18%
42%
50 %
Iterative Refinement (w G+E+P)
36%
50%
95%
7 %
9%
92%
8%
44%
42%
Table 9. Leveraging execution feedback helps reduce errors: Here we present the percentage of problems where the LM-generated
Kernel is for iterative refinement. We leveraging execution feedback helps the model achieve better correctness fasto, which is the percentage of problems where the model has at least one correct generation up to turn N 10. We note the various iterative refinement configurations, leveraging previous Generation G. Execution Result and Timing Profiles P.
F. Few Shot Experiment
For this experiment, we provide in-context examples of optimization techniques such as fusion, tiling, recompute, and asynchrony to models during kernel generation. As described Section 5.2.1, we provide three in-context examples: a fused GELU (Hendrycks Gimpel, 2023), a tiled matrix multiplication (Mills, 2024), and minimal Flash-Attention (Dao al., 2022; Kim, 2024) demonstrating effective shared memory I/O management. The prompt used for this experiment described in Appendix C.4. The speedup of these kernels computed over PyTorch Eager. We compare the performance of these few-shot kernels over the one-shot baseline below.
28

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Baseline
Few-Shot
Model
Level
fast
fasto
Kernel Length (chars)
fast
fast
Kernel Length (chars)
3%
27%
301018
6%
27%
360212
Llama 3.1-70B
0%
0%
646403
0%
0%
566668
0%
0%
404596
0%
4%
485332
1
10%
55%
343995
6%
39%
437768
OpenAI ol
2
24% %
56%
381474
16%
39%
432800
3
12%
56%
260273
8%
22%
364551
Table 10. Comparison of the Section 4.1 baseline and few-shot prompting performance across models. We examine the fast,, fast and cumulative character length of generated kernels per level.
77% of matrix multiplication problems in Level achieves speedup over the one-shot baseline through tiling. The runtime comparison for each GEMM variant is presented below as Table F.
Problem Name
Baseline (ms)
Few-Shot (ms)
Ref Torch (ms)
3D Tensor Matrix Multiplication
20.9
7.71
1.45
Matmul for Upper-Triangular Matrices
14
5.39
2.98
Matrix Scalar Multiplication
1.19
0.811
0,822
Standard Matrix Multiplication Matmul with Transposed Both Matmul with Transposed A 4D Tensor Matrix Multiplication Tall Skinny Matrix Multiplication Matmul with Diagonal Matrices
3.39
2.46
0.397
3.44
2.67
0.412
3.61
2.99
0.384
366
338
36
3.39
3.59
1.9
0.221
0.237
2.83
Table 11. Performance comparison of the Section 4.1 baseline and few-shot prompting level matrix multiplication problems.
Few-shot kernels generated for the following problems in level outperformed PyTorch Eager through aggressive shared memory I/O management.
Problem Name
Baseline (ms)
Few-Shot (ms)
Ref Torch (ms)
Cony2d InstanceNorm Divide
0.514
0.0823
0.0898
Gemm GroupNorm Swish Multiply Swish
0.124
0,0542
0.0891
Matmul Min Subtract
0.0651
0,0342
0.0397
Matmul GroupNorm LeakyReI.U Sum
0.0935
0.0504
0.072
Conv Transpoosd Swish GroupNorm HardSwish
33.3
29.6
35.2
Conv Tonnspsspps Mish Add Hardtanh Scaling
0.235
0.209
0.243
Conv Transpose3d Add HardSwish
15.6
14.1
22.2
Conv Tranppsssd Add Min GELU Multiply
0.365
0.349
0.4
ConnTrosspsee2 BiasAdd Clamp Scaling Clamp... Conv2d GroupNorm Tanh HardSwish ResidualAdd... Conv2d ReLU HardSwish
0.3
0.31
0.368
0.124
0.129
0.154
0,0681
0.0711
0.0768
Table 12, Performance comparison of the Section 4.1 baseline and few-shot prompting level for problems whose few-shot kernels outperform PyTorch Eager.
29

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
G. Cross-Hardwand Case Study
G.I. Evaluation across different hardware
To evaluate how generated kernels fare across different hardware platforms, we utilize a number of different NVIDIA GPUs that span different archtttrchi and capabilities. The specific details for each is provided Table 13.
Provider
GPU Type
Memory
Power
Microarchitecture
FP16 TFLOPS
Memory Bandwidth
Baremetal
NVIDIA L40S
GB
300W
Ada
362.05
864 GB/s
Baremetal
NVIDIA H100
80 GB
700W
Hopper
989.5
3350 GB/s
Serverless
NVIDIA L40S
48 GB
350W
Ada
362.05
864 GB/s
Serverless
NVIDIA A100
42 GB
400W
Ampere
312
1935 GB/s
Serverless
NVIDIA L4
24 GB
72W
Ada
121
300 GB/s
Serverless
NVIDIA T4
16 GB
70W
Turing
300 GB/s
Serverless
NVIDIA A10G
GB
300W
Ampere
125
600 GB/s
Table 13. Specifications of different GPUs, including memory, power consumption, micro-architecture,, FP16 TFLOPS, memory band- width, and their providers.
We ran the same set of kernels generated in Section 4.1 on variety of hardware (as listed in Table 13). We computed the fast, speedup against the PyTorch Eager baseline profiled on that particular hardware platform in Table 14.
Level GPUs
Llama-3.1-70b-Inst
DeepSeek-V? V3
DeepSeek-RI
L40S
3%
12%
H100
2%
7%
16%
A100
3%
7%
16%
L4
2%
4%
15%
T4
3%
7%
22%
A10G
2%
7%
12%
L40S
0%
4%
36%
H100
0%
4%
42%
A100
0%
4%
38%
L4
0%
4%
36%
T4
0%
4%
46%
A10G
0%
4%
47%
L40S
0%
8%
2%
H100
0%
10%
2%
A100
0%
8%
2%
L4
0%
6%
2%
T4
0%
10%
2%
A10G
0%
10%
0%
Table 14. KernelBench result across multiple hardware types: Speedup (fastı) over Torch Eager comparison of GPUs across different models levels. The kernels used across different GPUs are same as ones generated Single Attempt without hardware/plat- form specific information.
Based on the increased variability in fast, score for DeepSeek R1 as described in Section 4.4 and Table 14, we plot the individual speedups for each problem (in Levels and 2) across different GPUs. Speedup computed against Py Torch Eager and there a horizontal line at y = 1.0 mark the cutoff for fast,.
30

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Speedup per Probiem for Level 1: DeepSook-RR R1 against PoToch Eager (Correct Problems Only)
y-1.
LAGS
M1000
A.00
Speco.p
Problem
Figure 8. Speedup comparison across diffcrent GPUs for DeepSeck RI on Level (log scale).
Speedup per Problem Level 2: DeepSeekRR1 against PyTorch Eager (Correc: Problems Only)
LAOS
L4
A10G
Problem
Figure 9, Speedup comparison across different GPUs for DeepSeck-R1 on Level (log scale).
31

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
G.2. Effect of Providing Hardware Information
We provided sarewfie-specific information such as the GPU's micro-architecture, TFLOPS, and memory bandwidth (as shown in Table 13) to models during kernel generation. These generations were then profiled their respective hardware platforms. The speedup in terms of fast over PyTorch Eager is provided below in Table 15.
Level
GPUs
Llama-3.1-70b-Inst
DeepSeek-V3
DeepSeek-RI
OpenAI-ol
L40S
2%
5%
14%
13%
1
H100
3%
6%
14%
9%
A100
4%
4%
16%
9%
L40S
1%
4%
33%
25%
2
H100
2%
2%
35%
33%
A100
2%
1%
31%
32%
L40S
0%
2%
0%
10%
3
H100
2%
8%
0%
12%
A100
0%
0%
2%
8%
Table 15. Providing hardware specifications in context: Specdup (fastz) over Torch Eager comparison of GPUs across different models and levels.
As described in Section 5.2.2, DeepSeek-Ri R1 begins to generate kernels with warp matrix instructions. We show the generated kernel for Problem Level I when provided with H100 specifications in context below.
(wmma)
Vinclude <torch/extension h>
Vinclude <cuda_fp16. f.f6.p11... h>>
#include <cuda_runtime h>>
#inolude <mma.ht
6
using namespace nvcuda;
global void hgemm kernel
hgenn_kernel (const half* &, const halfs B, half* c, int N)
const int WARPS. PERPPEERR__ BLOCK 8;
10
const int WMMA_P 16;
11
const int WMMA_N N 16;
12
const WMMA_K 16;
13
14
shared_.. half Ashared [ WARPS PER PLOCK) [MMA M]
[MMA K.
15
shared half Bshared WARPS PER BLOCK] [MMA X]
[MMA N)
16
17
int warpId threadIdx.x 32;
18
int lancId Id threadIdx. 32;
19
20
int tileM (blockIdx.y WARPS_PPS
warpId) WMMM_M;
21
int tileN blockIdx.x WMMA_N;
22
23
wmma: : f www.mentgttt......mm.......... <wmma: :atrii a, WXMA M, WMMA N, NMMA K, half, w wwtpw_mjjor>.. a_frag;
24
wmma: aaagenn<<wmaa :matrii WAKA M, MMMA N, WMMA_ half, w b_fragi
25
wmma: Fragment<wmma httpmeenttwmmm::accmmllllaaaaaa taccumllator, WMMA_M, WMMA_N, WMMA_K, half> c_frag; wmma: fill fralm_fragment/ (c_frag, 1
26
27
28
for (int k = 0; k < N; k += WMMA K) 1 int loadAidx = laneId (tileK + (laneId / WMMA_ K) N k (LaneId $ MMA K) int loadBidx laneId (k (laneId WMMA N) N tileN (LaneId $ WMAA N)
29
30
32
if (tileM (laneId / WMMA_K) N & + (LancId WMMA_K) < N) Ashared[warpIdl (laneId / WMMA_] [lancId $ K] = A dxxa else
33
34
35
Ashared (varpid)] [laneld / шмMA_] [laneld WMMA_K] Float2hal (0.0f))
36
32

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Figure 10. A CUDA kemnel generated by DeepScek-RI I for Level Prohlem when provided with hardware-speciffc information the
H100 GPU.
37
if (k + (laneId / WMMA N) N & tileN + (laneId WMMA N) < N)
38
[waprid[ [laneId / WMKA_N] [laneld $ WMMA_N] - B [loadBidx]
39
else
40
Bshared[warpid. [laneId WM [laneId WMMA Float2half (0.
42
syncthreads (
43
44
-Lcad sync (a_ frag, 6Ashared wwrpId] [0] [0], WMMA K)
wmma: wwwdd syne frag,
&Bshared [warpId] [0] [0), NMMA_N)
46
wnmaa: : :ww.a syncaassscccc..cccm (c_ frag, trag, b_ orag, c_trag);
47
48
syncthreads ())
49
50
51
int storeCidx (tileM + (laneId / WMA [)) + N tileN (laneId 8 WMMA N)
52
if (tileM (laneId / WMA N) N & tileN + (laneId WMMA N) < N)
53
C[storeCidx] e fffra..x(aanntl
54
55
56
torch: Tenscr hgemm_cuda (torch:: :Tensor A, torch: Tensor B, int N) { auto C torch: ([N, N), A.eptions () dttyp :otorcc::kFoattttoo )
57
58
const int WARPS PER_EEEBLOOO BLOCK 8; dim3 (IN + 15) / 16, (N + 15) / (16 + SSRP__PER_BLOCS )
60
61
dim3 b3ock(32 WARPS PER BLOCK)
62
63
block>>> (A. M..data (),, B.data_ (), c.data_ptr<half
4 (), N);
64
return Cj
65
G.3. In-context examples applying hardware-specific instructions
For this experiment, combine the approaches of Sections 5.2.1 and 5.2.2 by providing in-context examples using architecture- sppecic specific instructions such (1) wmma and (2) memepy.async along with the same hardware-specific information in Table 13. We specifically target 100 GPUs since the Ampere architecture supports both Tensor Core operations and Asynchronous Memory Movement. We use DeepSeek-RI based on its ability to generate kernels with wmma instructions without explicit examples as shown in Section 5.2.2 (though fails to utilize them correctly). Finally, we target simple KernelBench matrix multiplication problems (17 of the Level problems).
Qualitatively, providing DeepSeek-RR R1 with in-context example using WMMA instructions motivated the model to try to apply WMMAs as much as possible (all matrix multiplication problems we were targeting had generated kernels that included WMMAs). However, only out of those 17 kernels were correct. The kernel that achieved a >> 3.5x speedup over Py PyTorch Eager performed a matrix multiplication for two upper triangular matrices and skipped unnecessary computation, an optimization enabled by prior knowledge about input matrix characteristies rather than better Tensor Core utilization. The remaining 4 kernels were slower than PyTone Po'ch Eager.
33

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Speedup per Problem for Level 1 MatMul Problems: DeepSeek-R1 against PyTorch Eager on A100 (WMMA)
3.5
3.0
2.5
Speedup
2.0
1.5
1.0
0.5
0.0
11
12
13
74
15
16
17
18
Problem ID
Figure 11. Speedup over PyTorch Fager when providing DeepSeck-RI with an in-context example using WMMA instructions perform basic matrix multiplication.
On the other hand, providing DeepSeek-RI with in-context example using memepy _asyne instructions also motivated apply these instructions to all of its generated kernels. However, none of the kernels were correct. The fact that DeepSeek-RI points certain hardware capabilities being easier for DeepSeek-RI to exploit than others.
H. High-Throughput Evaluation System
H.1. Single-shot Experiments: Batched Kernel Generation
Given the high volume of GPU kernels to evaluate, we build a fast and highly-parallelized evaluation system, where we separate into the kernel generation and evaluation process into 3 stages, as shown in Figure
Inference: We query LMs in parallel and store the generated kernel.
CPU Pre-Compile: We compile the model-generated kernels with nvec for a specified hardware into a binary, parallelized on CPUs and each kernel binary is saved to their individual specific directory for caching.
GPU Evaluation: With the kernel binary already built on CPU, focus evaluating multiple kernels in parallel across multiple GPU devices. However, ensure accurate kernel timing, only evaluate one kernel at time device.
H.2. Iterative Refinement Experiments: GPU Orchestrator System
Based on the single-shot system, we also design a platform to handle multiple iterative refinement experiments at once. We treat each iterative refinement experiment as a finite state machine, where the states are L.M-daedd kernel generation, pre-compilation, kernel execution, and profiling. The transitions are based on environment feedback, and can change based on different experiment setups.
34

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Inference
Pre-Compile
GPU Eval
Thread Paraile! (-100) over API requests
Device over GPUs (max 8)
over CPUs (-100)
Datasets
Fetch
Eval Proc.
Fetch Code
Problem
Prompt Construction
Code
Evall Proc.
Isolation
Thread Pool
Eval Proc.
Query API
Query API
Compile cuda binaries
Cache hit
Eval Proc.
Response Parsing
Cache Dit/run/ppoolemmmmmmmpeehhaaat
Log Resulta
Figure KernelBench provide a high throughput kernel generation and evaluation system. We parallelized generation, compilation. and evaluation of kernels across CPUs and GPUs.
Our system was run on a node with available GPUs. Unlike the single-shot system, batching each generation and kernel execution is highly inefficient thus, we design pipelined, multiprocessing system with a GPU orchestrator with the following characteristics:
The orchestrator spawns multiple independent processes that each handle an independent task in KernelBench. These processes run the multi-turn state machine logic for the iterative refinement experiments only the kernel execution state requires acquiring a GPU.
CPU Parallelism:
The GPU orchestrator keeps separate process running that handles which processes acquire GPU using semaphores. Processes can request GPU from this process when it is ready to execute and evaluate kernel code. We try to minimize process control over GPU to maximize resource throughput, given system with a limited number of available GPUs.
Acquiring GPUs:
Pre-compiling the CPU: To avoid processes hogging GPU time, we pre-compile kernels with nvec the CPU for specified hardware into binary. We also did this same trick for the single-shot system, but for separate reasons.
Evaluating Kernels on the GPU: and profiling. We found that waiting on GPUs is the primary bottleneck in the orchestrator, so we designed the orchestrator maximize device occupancy.
The only state where the finite state machine uses the GPU is for kernel execution
The system generally supports overlapping the generation of kernel code and the execution of erready-generated kernel code. There are also several unavoidable errors such as CUDA illegal memory accesses and deadlocks due to faulty kernel generations that the orchestrator solves by releasing and spawning new processes when encountered, and we wrote specifically handlers to ensure these errors are properly captured without crashing the orchestrator itself.
H.3. UI: Visualizing Kernel Generation Trajectories
To qualitatively observe the generated and compare them across techniques, we design an interface to easily visualize them. provide this as part of the KernelBench framework.
35

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Evaluation Result for Sample
Completion:
Seccess
Cormetneess
Pass
Runtime: 4.17 #s
Booel ne Torth Time: 4.22
Breeline Torch Complee Time: 2:63
Metadera::
La "udle?" 4,39, 4,34, 0.0085e, 'mcan's
"comallee's Trux,
"erg.rnniaee "rantinc alels: trtelal: 571))
Teudles
"handwere"
"IVESAA LABS".
Reference Torch Code
Generated Kernel Sample
Num Source Tokens 592
Nam Tokens:
import to
import to rchunn as an import tarch.nn. functional
terch
Laport terch.nn, fros torch.utils, epp_erceniion Inport Intine
Mode2Inn.Modutell expansion
CUDA for elementuise additiin lementsisee source
def
init Iself, in channels, chameels, stride:1::
...
#include crarch/esstesiinn.b #include =cudi__unttiec... runtiec.lo-
:paran iichhmnlaa: Naddur impul chamnls
:paran out, cchnnels:: Number of ouiput channets
:paran stride: Stride for the first convolutional layer Downsarple
globbl___ elanentwise, add kei.eeeeeee floate fleab- sizel
shortcut
int blockldx.x blocsDin.x threadids.xi lidx < sizel
superiMadel, self). init
cud[ide] a[ids] b[idk];
self.conee in www.aaaiinncceeeeeeeee out cha_nnh
Kerne
bias=False)
self.bnl www.attoennoettrtt.oooo..c
Figure 13. We provide visual interface for kernel inspection. This allows casily examine kernel content, its performance, and compare across various techniques and configurations.
I. KernelBench Score Design: Comparing Correctness and Performance
Evaluating kernel generation models requires balancing crucial aspects: correctness and performance. Correctness ensures that the generated kernel produces the expected outputs, while performance determines the efficiency of execution, typically measured in speedup. Since optimizing one aspect can sometimes degrade the other (e.g., aggressive optimizations may introduce more correctness errors), KernelBench uses the fast, metric. In this section, we also include some other metrics we have considered and compare them. Table 16 first shows correctness and speedup separately.
Level
Level
Level
Model
% Correct
Speedup
% Correct
Speedup
% Correct
Specdup
3.5-sonnet
53.0
0.329074
10.0
1.026136
12.0
0.448122
deepseek- V3
44.0
0.252212
5.0
1.023760
30.0
0.778776
deepseek-r1
67.0
0.483147
62.0
1.039562
8.0
0.685164
gpt-4o
39.0
0.330257
9.0
1.000646
14.0
0.612829
1lama-3.1-400
40.0
0.228077
0.0
0.000000
4.0
0.944302
Ilama-3.1-706
27.0
0.178835
0.0
0,000000
0.0
0.000000
openai-ol
55.0
0.325485
56.0
0.844315
56.0
0.783615
Table 16. Correctness and Speedup of Models across different levels. Speedup here measured the geometric of the correct samples only.
I.I. Metries Exploration
The various metrics we explored are defined follow.
Let & be an indicator variable denoting correctness of the i-th sample, where:
1, if the solution is correct
C 0, otherwise
Let To, be the baseline execution time and Tg the actual execution time for the generated kernel. The speedup ratio for
36

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
sample defined
Tb. S = T
Let be the total number of samples and Roooeett
Li=1 Cj be the number of correct samples.
Model Ranking Based on Each Metric
fast_p_0
(.0533
fast_p_1-
(0.10)
(0.04)
(0.03)
(0.03)
arithemtic mean speeddssee____ee speed_ratio___eeeeeeeeeeeeeeetee rratrrattoootteee corlect_only (0.81)
(0.73)
1 (0.88)
6 (0.56)
(0.63)
(0.46)
(0.77)
arithemtic haatemttccmeansseee speca___pee__ati
(0.43)
(0.32)
1 (0.59)
6 (0.22)
5 (0.25)
(0.12)
3 (0.43)
arithemtic_noe sethemticmmean..oo mean, mean_speedd ratio lazy
(1.15)
(1.12)
1 (1.18)
6 (1.02)
(1.100
(1.02)
(1.15)
Metric
geometric mean speeed_______ee ratio ceorectttttoo only
(0.33)
(0.25)
(0.48)
2 (0.33)
(0.23)
(0.18)
(0.33)
geometric_men sacls__meaanaaaeeeeee __e___mmeeedd_____ee__
55)
(0.55)
(0.63)
geometric_ment meean_eeaaa__eee
_ttp_maeeedaaaaaaeeeeaamae sepessppeeeiiie ratio_layyy
(1.04)
(1.04)
1 (1.07)
(1.01)
(1.03)
(1.01)
(1.05)
geomeerii mean, mean, mean_speed ratio
(0.17)
(0.11)
1 (0.32)
4 (0.13)
(0.09)
(0.05)
(0.18)
adjustable weighted_ score
3 (47.89)
4 (38.07)
(62.89)
(30.37)
(32.55)
7 (19.69)
(48.79)
weighted_log locscceeeee
2 (23.89)
4 (17.81)
1 (34.28)
(15.68)
(13.87)
(8.81)
(23.89)
claudee 2.5 S-somet
deeponek-V33 v/
deepsack-r1
got-4o
Bama- 2.1 1-405b
lama- 3.1 1-770b
openal-g g1
Model
Figure 14. Model Performance Ranking Across Various Scoring Metrics
Figure 14 compares model rankings under different performance metrics, capturing variations in correctness and execution speed. Each cell displays the ranking and the corresponding raw score for a given model and metric. Differences in rankings highlight the impact metric design on performance evaluation.
Arithmetic Mean Speed Ratio (Correct Only)
AMSR
S
Flcorrect 1=1
Arithmetic Mean Speed Ratio (All Samples)
1 Σ
AMSR
c;S;
1=1
Arithmetic Mean Speed Ratio (Lazy)
1
AMSRay
max(c,S;,11
im
37

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Geometric Mean Speed Ratio (Correct Only)
72
Pouned
GMSR correct
II
Geometric Mean Speed Ratio (All Samples)
GMSR
II [=1
S
Geometric Mean Speed Ratio (Lazy)
GMSRlag
m 1)
Scaled Geometric Mean Speed Ratio
Tocorrot
SGMSR
GMSR correct
n
Adjustable Weighted Score
AWS Σ c (correctness weight + perf weight x S:)
where correctness and performance weights satisfy:
correctness weight + perf weight
We use correctness heggehesss.wighh--ef...eehhhhhhe in our experiments.
Weighted Log Score
WLS Σ c log(1 + S.) =1
38

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
J. KernelBench vs. LiveCodeBench Correlation
To assess the relationship between general coding ability and kernel-specific performance, we compare model rankings across LiveCodeBench (Jain al., 2024) and KemelBench (Levels 1-3).
Heatmap of Rankings
gpt-4o
(32.9)
(0.04)
(0.05)
(0.0)
openai-o1
(71.0)
2 (0.1)
(0.24)
1 (0.12)
deepseek-V3
(34.6)
(0.06)
(0.04)
2 (0.08)
deepseek-r1
(62.8)
(0.12)
(0.36)
(0.02)
daude-3.5-somee
(37.2)
(0.07)
(0.02)
lama-3.1-70b
(26.0)
(0.0)
laama3.1-405b
LiveCodeBench
KernelBench_Lvee KernelBench_Levee KermelBench_Lever
Figure 15. Heatmap of model rankings across LiveCodeBench and KemelBench (Levels 1-3).
observe that LivcCodeBench performance tend to correlate strongly KernelBench, though the rankings are not perfectly aligned. 1lama-3.1-70b and Ilama-3.1 -400bb consistently rank the worst across benchmarks, suggesting that weaker general coding ability correlates with poor kernel-specific performance. Openai-ol, which ranks 1st in LiveCodeBench, remains highly competitive in KernelBench (2nd in Level I and 2, 1st in Level 3). Deepseek-rl, ranking 2nd in LiveCodeBench, achieves 1st place in KernelBench Level and Level 2.
While models that perform well in LiveCodeBench generally achieve strong results in KernelBench, the variability in rankings across different levels of KernelBench suggests that additional skills are required for high performance in kernel-specific tasks. For example, deepseek-rl, which ranks 2nd LiveCodeBench, drops to 4th in KernelBench level 3, indicating that some aspects of kernel optimization may favor different model strengths.
K. KernelBench Tasks Breakdown
K.1. KernelBench Tasks
Reference KernelBench tasks are given in FP32, and given a tolerance threshold (1e-2), using lower precision solutions is valid. Furthermore, the problem size of each task is fixed, since our goal is specialized fast kernels not generally fast kernels for any arbitrary shape.
Here we list the names of the KernelBench tasks categorized by level.
39

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
List of Tasks of All Levels
Level Task Names
1. Square matrix multiplication
35. GroupNorm
2. Standard matrix multiplication
36. RMSNorm
3. Batched matrix multiplication
37. FrobeniusNom
4. Matrix vector multiplication
38. LINorm
39. L2Norm
5. Matrix scalar multiplication
40. LayerNorm
6. Matmul with large K dimension
41. Max Pooling ID
7. Matmul with small K dimension
42. Max Pooling 2D
8. Matmul with irregular shapes
43. Max Pooling 3D
9. Tall skinny matrix multiplication
44. Average Pooling ID
tensor matrix multiplication
45. Average Pooling 2D
11. 4D tensor matrix multiplication
46. Average Pooling 3D
12. Matmul with diagonal matrices
47. Sum reduction over dimension
13. Maimul for symmetric matrices
48. Mean reduction over dimension
14. Matmul for upper triangular matrices
49. Max reduction over dimension
15. Matmul for lower triangular matrices
50. Product reduction over dimension
Matmul with transposed
51. Argmax over dimension
17. Matmul with tramsposed B
52. Argmin over dimension
18. Matmul with transposed both
53. Min reduction over a dimension
19. ReLU
54. conv standard 3D square input square kernel
20. LeakyReLU
55. conv standard 2D asymmetrie imput square kernel
21. Sigmoid
56. conv standard 2D asymmetric imput asymmetric kernel
57. conv transposed 2D square input square kernel
Tanh
conv transposed asymmetric input asymmetric kernel
23. Softmax
59. conv standard 3D asymmetric input square kernel
24. LogSoftmax
60. conv standard 3D square input asymmetric kernel
25. Swish
61. conv transposed 3D square input square kernel
26. GELU
62. conv standard 2D square input asymmetric kernel
27. SELU
63. conv standard 2D square input square kernel
HardSigmoid
64. conv transposed
29. Softplus
65. conv transposed 2D square input asymmetric kernel
30. Softsign
66. conv standard 3D asymmetric input asymmetric kernel
31. ELU
67. conv standard 1D
32. HardTanh
68. conv transposed square input asymmetric kernel
33. BatchNorm
69. conv transposed 2D asymmetric input asymmetric kernel
InstanceNorm
70. conv transposed asymmetric input square kernel
40

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
71. conv transposed 2D asymmetric input square kerme
84. conv depthwise 2D asymmetric input square kernel
72. conv transposed 3D asymmetric input asymmetric kemne strided padded grouped
85. conv depthwise 2D asymmetric input asymmetric kernel
86. conv depthwise separable 2D
73. conv transposed 3D asymmetric input square kernel strided padded grouped
87. conv pointwise 2D
74. cony transposed ID dilated
88. MinGPTNewGelu
conv transposed 2D asymmetric input asymmetric kernel
89. cumsum
strided grouped padded dilated
90. cumprod
76. standard ID dilated strided
91. cumsum reverse
77. conv transposed 3D square input square kernel padded di- lated strided
92. cumsum exclusive
78. conv transposed 2D asymmetric input asymmetric kernel padded
93. masked cumsum
94. MSELoss
cony transposed ID asymmetric input square kernel padded
strided dilated
95. CrossEntropyLoss
80. conv standard 2D square input asymmetric kernel dilated padded
96. HuberLoss
97. CosineSimillrityL.os
81. transposed 2D asymmetric input square kernel dilated padded strided
98. KLDivLoss
82. depthwise 2D square input square kernel
99. TripletMarginLoss
83. conv depthwise 2D square input asymmetric kernel
100. HingeLoss
Level Task Names
1. Conv2D ReLU BiasAdd
18. Matmul Sum Max AvgPool LogSumExp LogSumExp
2. Conv ConvTranspose2d BiasAdd Clamp Scaling Clamp Divide
19. ConVTranspos22d GELU GroupNorm
3. Con/Transpose3d Sum LayerNorm AvgPool GELU
20. ConvTranspose3d Sum ResidualAdd Multiply ResidualAdd
4. Conv2d Mish Mish
21. Conv2d Add Scale Sigmoid GroupNorm
22. Matmul Scale ResidualAdd Clamp LogSumExp Mish
5. Conv ConyTranspose2d Subtract Tanh
23. Conv3d GroupNorm Mean
6. Conv3d Softmax MaxPool MaxPool
24. Conv3d Min Softmax
7. Conv3d ReLU LeakyReI.U GELU Sigmoid BiasAdd
25. Conv2d Tanh Tanh
8. Conv3d Divide Max AvoPalAvgPPol BiasAdd Sum
26. ConvTranspose3d Add HardSwish
9. Matmul Subtract Multiply ReLU
27. Conv3d HardSwish ReLU Softmax Mean
10. Con/Transpose2d MaxPool Hardtanh Mean Tanh
28. BMM InstanceNorm Sum ResidualAdd Multiply
Conv Tranppsasp BatchNorm Tanh MaxPool GroupNorm
29. Matmul Mish Mish
12. Gemm Multiply LeakyReLU
30. Gemm GroupNorm Hardtanh
13. Con/Transpose3d Mean Add Softmax Tanh Scaling
31. Conv2d Add Multiply
14. Gemm Divide Sum Scaling
32. Conv2d Scaling Min
15. ConvTransposeee BatchNorm Subtract
33. Gemm Scale BatchNorm
16. Conv ConnTraaspose2d Mish Add Hardtanh Scaling
34. ConvTranspose3d LayerNorm GELU Scaling
Conv2d InstanceNorm Divide
35. Conv2d Subtract HardSwish MaxPool Mish

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
36. Conv ConnTraaspose2d Min Sum GELU Add
68. Matmul Min Subtract
37. Matmul Swish Sum GroupNorm
69. Conv2d HardSwish ReLU
38. ConTTranspoccdd AvgPool Clamp Softmax Multiply
70. Gemm Sigmoid Scaling ResidualAdd
39. Gemm Scale BatchNorm
71. Conv2d Divide LeakyReLU
40. Matmul Scaling ResidualAdd
72. ConyTranspose3d BatchNorm AvgPool AvgPool
41. Gemm BatchNorm GELU GroupNorm Mean ReLU
73. Conv2d BatchNorm Scaling
42. Cony Transsasspose2d GlobalAvgPool BiasAdd LogSumExp Sum Multiply
74. ConvTranspose3d LeakyReLU Multiply LeakyReLU Max
75. Gemm GroupNorm Min BiasAdd
43. Conv3d Max LogSumExp ReLU
76. Gemm Add ReLU
44. Conv Transraappoee2d Multiply AvgPool GlobalAvgPool
Mcan
77. ConnTransposedd Scale BatchNorm GlobalAvgPool
45. Gemm Sigmoid Sum Lxp
78. ConvTranspose3d Max Max Sum
46. Conv2d Subtract Tanh Subtract AvgPool
79. Conv3d Multiply InstanceNorm Clamp Multiply Max
47. Conv3d Mish Tanh
80. Gemm Max Subtract GELU
48. Conv3d Scaling Tanh Multiply Sigmoid
81. Gemm Swish Divide Clamp Tanh Clamp
49. Conv ConnTranspose3d Softmax Sigmoid
82. Conv2d Tanh Scaling BiasAdd Max
50. Conv Tonnrraaspoed3d Scaling AvgPool BiasAdd Scaling
83. Conv3d GroupNorm Min Clamp Dropout
Gemm Subtract GlobalAvgPool LogSumExp GELU Residu- alAdd
84. Gemm BatchNorm Scaling Softmax
85. Conv2d GroupNorm Scale MaxPool Clamp
Conv2d Activation BatchNorm
86. Matmul Divide GELU
53. Gemm Scaling Hardtanh GELU
87. Conv2d Subtract Subtract Mish
54. Conv2d Multiply Le GELU
88. Gemm GroupNorm Swish Multiply Swish
Matmul MaxPool Sum Scale
89. Conv ConvTransposeia MaxPool Softmax Subtract Swish Max
56. Matmul Sigmoid Sum
90. Conv3d LeakyReLU Sum Clamp GELU
57. Conv2d ReLU HardSwish
91. Con'Transpose2d Softmax BiasAdd Scaling Sigmoid
58. Conv Transposed LogSumExp HardSwish Subtract Clamp Max
92. Conv2d GroupNorm Tanh HardSwish ResidualAdd LogSum- Exp
Matmul Swish Scaling
60. Conv Transpssspoed Swish GroupNorm HardSwish
93. Tranppsspp Add Min GELU Multiply
61.
ConnTraaspocedd ReLU GroupNorm
94. Gemm BiasAdd Hardtanh Mish GroupNorm
62. Matmul GroupNorm LeakyReLUU Sum
95. Matmul Add Swish Tanh GELU Hardtanh
Gemm ReLU Divide
96. ConvTranspose3d Multiply Max GlobalAvgPool Clamp
64. Gemm LogSumExp Lea LeakyReLU GELU GELU
97. Matmul BatchNorm BiasAdd Divide Swish
65. Conv2d AvgPool Sigmoid Sum
98. Matmul AvgPool GELU Scale Max
66. Matmul Dropout Mean Softmax
99. Matmul GELU Softmax
67. Conv2d GELU GlobalAvgPool
100. ConTTranspose3d Clamp Min Divide
Level 3 Task Names
42

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
1. MLP
18. SqucezeNet
35. LTSM
2. ShdeLowWideMI.P
19. MobileNetVI
36. LTSMHn
3. DeepNarrow.MLP
20. MobileNetV2
37. LTSMCn
4. LcNet5
21. FificientNerMMRConn
38. LTSMBidirectional
5. AlexNet
22. EfficientNeiB0
39. GRU
6. GongleNeetieeetiioMMouulle
23. EfficientNetB1
40. GRUHidden
7. GowegleNelineptioonVi
24. EfficientNetB2
41. GRUBirectional
8. ResNetBasicBlock
25. ShuffleNetUni
42. GRUBiddircciionallHddee
9. ResNet18
26. ShuffleNet
43. MinGPTCausslAAttention
10. ResNet101
27. RegNet
44. MiniGPTBlock
11. VGG16
28. VisionTransfon Transformer
45. UNetSoftmax
VGG19
29. SwinMI.P
46. NetVHaaWWihCChhoCCusseer
13. DenseNedd22Transitioool
30. SwinTraansfrmerr22
14. DenseNor121DeeseeBock
31. VisionAttention
47. NetVlalNoGhossChastes
DenseNet121
32. Convolutional Transformer
48. Mamba2RetumY
16. DenseNet201
33. VanillaRNN
49. Mamba2RetuuminnlSStae
17. SquccraNatFrrMMddule
34. VamillaRNNHidden
50. ReL.USeffAttentio
K.2. Level 2 Synthetic Generation
want especially elaborate on the design and construction Level problems. The construction of level done by randomly picking one main operator and to epilogue operators, for a total of 3 to 6 operators per task. Figure 16, which highlights the relative frequency of different task sizes.
The mainloop operators include Matmul, BMM, Conv2d, Conv3d, Con/Transpose2d, and CownTTransoseeaee. as shown in Figure 17.
The epilogue operators divided into different classes:
activations:
ReLU, Sigmoid, Tanh, LeakyReLU, GELU, Swish, Softmax, Mish, Hardtanh, WardSwish element-wise ops: Add, Multiply, Subtract, Divide, Clamp, Scale, ResidualAdd
normalizations: BatchNorm, LayerNorm, InstanceNorm, GroupNorm
pooling: MaxPool, AvgPool, GlobalAvgPool
bias: BiasAdd
reductions: Sum, Mean, Max, Min, LogSumExp
others: Dropout, ResidualAdd, Scaling
The distribution of epilogue operators in Level 2 tasks is illustrated in Figure 18.
L. Kernel Fusion Investigation
Kernel Fusion an important optimization deployed optimizing Deep Learning programs. By fusing multiple operations into single kernel, the optimized program can reduce memory traffic which improves performance especially as data movement is expensive. Just to reiterate, in KernelBench the model has full flexibility to decide what subset of operators in the PyTorch reference to optimize and fuse. We believe this is one of the crucial abilities when model is given distinct or new architectures in the real-world setting.
KernelBench's s 3-level categorization helps disentangle fusion decisions and kernel generation. Level problems (single operators) only test the model's ability to write optimized kernels; Level 2 and 3 problems are designed to additionally evaluate the model's ' ability to identify and leverage fusion opportunities.
43

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Histogram Number Operators in Level Tasks
30
25
Requeecy
20
15
10
3.0
6.5
7.0
Number Operators
Figure Histogram the number operators per task Level
Ilssogram Mein Operators Level Teoks
Coliegue Tasks
PODEANY 13
20-
coneaa
Ber
Narmoall
Mith Aaniiban
Open
Figure 17. Histogram of main operators in Level 2 tasks.
Figure 18. Histogram of epilogue operators in Level 2 tasks.
From the baseline results in Section 4.1, we investigate the particular fusion choices made by the language model and how these influence the runtime of these kerelss. We focus specifically on Level tasks, where fusions can be applied to all problems. They are composed with one mainloop (e.g. conv, matmul) and 2 epilogue operations (non-linearities reductions, etc). Figure 17 and 18 shows distribution of these operators in the 100 Level programs.
We focus analysis on the behavior of DeepSeek-R1. R1, which has the best performance Level 2 problems shown in Table 1.
L.1. LM-generated Kernel Fusions Level 2
We manually inspected all 100 generated programs and note down the fusion patterns. Model always attempts generate 1-2 fused kernels per problem: 88 problems have I fused kernels, and 12 problems contain 2 fused kernels.
One interesting point of analysis is how aggressively the LM chooses to fuse operators. In Figure 19, we plot a histogram to inspect the size (in terms of ratio of total operators) of generated fused kernels Level 2. Specifically, for each kernel we count number of operators fused compared to the total number operators in the program. For example, suppose program has operators, and the LM chooses to fuse the first two and last four operators. Then we add 33.3% and 66.6% to our list.
44

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Distribution Fusion Ratios Buckets
27
18
Count Problems
17
15
5
0-10%
10.20%
20 30%
30-40%
40-50%
50 60%60%
60-70%
70-80%
80 80-90%
90-100%
Ratio Fused Operators Total Operators
Figure 19. Distribution of Fusions by DeepSeek-RI on Level Problems. We bucket (in increments of 10%) all fusions performed by DeepSeck-RI on Lovel based on the percentage operators it fused in the full task. We inspect the fused kernels terms of ratio of total operators of that program. Specifically, for each kernel we count number of operators fused compared to the total number of operators in the program.
We observe that the models attempl to fuse more than half the operators on average. Only 18% of programs fuse all operators program into a single kernel.
L.2. Understanding Fusion in Slower Kernels
To understand the quality of the fusion decision and how poor fusion leads to low performance, we analyzed the DeepSeek-Ri Level 2 kernels that were slower than Py Torch Eager (as shown in Table 17) and drew two observations:
1. Main loop operators (e.g., Conv) were not fused with epilogue operators.
2. The model's attempt to fuse main loop operators (e.g., GEMM other ops) was faster than launching highly optimized CuBLAS kernels.
L.3. Comparison with Auto-Fusion techniques
We focus comparison on widely-adopted torch. compile le in PyTorch 2 (Ansel et al., 2024) which employs an auto- fusion policy over TorchInductor's de-ine-by-run IR. In Table 17, we can see the difference in fusion decisions of R1 and torch.compile in Table XX. torch. compile often creates sophisticated fusion patterns, by breaking Convolutions or GroupNorm into smaller multi-pase passs kernels that compute parietal results and statistics in parallel, something RI Rl-kemelss rarely do with room for improvement. However, 37% of RI-kernels actually are faster than torch. compile, as torch. comp i le has to fit pre-existing defined templates, whereas RI can generate more custom optimizations.
45

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Problem
Speedup
Sequence of Operators
Total
LLM Generated
torch. compile
0.1810
[geman, relu, divide]
Kgemm, relu, divide)| (matmul, swish, scaling)]
[gemm, (relu.div)
59
0.1996
Imatmul, swish, scaling]
[gemm, (sigmoid.
35
0.2331
Jcoav2d, subtract. hardswish, maxpool, mish]
[conv2d, (subtract, hardswish), (maxppol, mish)]
(convolution, convolution, convolution,
(convolution, auin suh.hanlswish),
(convolution, .uuh
bardswish,mas wtwpshhw...oooooeea. pow.oooll.dlll.. with wnth.iiieechiceiteeiiie ind.iiidddcsiiill
27
0.5501
Jeeav3d, hardswish, relu, softmax, mean|
[conv, (bardswish, relu), (seltmax, mean)]
[convolution.
(convolution, hurdiinnn,nndwwie hu stl.
(convolu-
tiom wom.hankswish.rele softmax,,menn
0.6156
[gemm, max, subtract, gelu]
[gemm,, (max, subtract, gelu)] [mmatmul, swish, sum, groupoorm]
[gemm, (max, (ew.eannuu.gglli)
37
0.6922
[matmul, swish, sum, groupnorm]
[gemm, native group.norm,
native group aorm]
54
0.7394
[coav2d, multiply, leakyrelu, gelul
4
[cons, (maltiply, leakyrelu, gelu)]
[convolation, convolution, convolation, (convolution, mul. mual/eaky..rr mul.leskly.rluugelu connvllution, convolation, convolation, wwwp:o/tiioo.nnowss...ee..
69
0.7433
Jccav2d, 2d, hardswish, relu)
[cons, (hardswish, relu)]
0.7701
[conv2d, min, tanh, tanh]
[cons, (min, tanh, tanh)]
[convolution, convolation, convolution. (convolution.min.tal))
0.7859
Isonvtranspose2d, min, sum, gelu, add]
tconvtranspose, (min, sum, gelu, add))
connvolution, convolution, convolution, (convolution.ming min), sum, (gelu,ad))
0.8231
[conv2d, divide, leakyrelu]
[com22d, (divide, leakyrelo))]
[convolution, convolution, convolation, detovolution.di.leeaal [rlll]
32
0.8399
[conv2d, scaling, min]
[cons, (scaling, min)]
[convolation, convolation, convolution, (convolution.n) mul, mini]
68
0.8411
[matmul, min, subtract]
[matmul. (min, subtract)
[gomm. (minimum,subb) Igemm, (logsumexp pea.y..aaaaaayapaka relu.eelll)))
0.8545
Igemm, logsumexp. leakyrelu, leakyrelu. gelu, gelu] [sonv2d, avgpool, sigmoid, sum]
Igemm, logsumexp. (leakyrelu, leakyrelu, gelu, geluj]
0.8615
[cons, (avgpool, sigmoid, sum)]
[convolation, convolation, convolution,
convolation.
(convolutio,,agg pool22,sigmoiddsm))
0.8622
[matmul, divide, gelu] [ceav2d, subtract, tanh, subtract, avgpool)
[matmul, (divide, gelo)]
[gemm, (div.gelo)
46
0.8671
(conv2d), (subtract, tanb, subtract). (avgponl))
[convolution, convolution, convolution,
(convolution. sub.ianh),
(convolution, www.tuhh..vgg...olll.l.
0.8715
[gemm, multiply, leakyrelu]
[gemm, (multiply, leakyrehu)
[gemm, mul.Jeakkyrrlul
87
0.8930
lecav2d, subtract. subtract, mish]
[cons, (subtract, subtract. mishi]
cconvolution, convolution, convolution,
cconvooutio.suubmiiii
0.9063
[sonv3d, mish, tanh]
[cons, (mish, tanh)]
[convolution, (comvolutioo.miss..aal)
82
0.9092
[coav2d, tunh, scaling. biasadd. max|
[conv2d, (lanb, scaling, biasadd), max]
convolutiion convolution, convolution,
(convolution. htt.://tttnnnmmllmllii.
max poo.ld_wwthhhppmoo [iiiii_iittii
0,9380
[conv2d, softplus, tanh, multiply. batchnorm]
[conv2d, (softplus, tanh, multiply), batchnorm]
[convolution, convolution, convolution. wwt.ortttioooooolssoooo canlllssttn.mmm.
_native batch.norm. legit. functional), fon.etig....
(convolution. sotpooossosopoooollloooo .aaive www.hh.nooo........... .nw.o.llleee finnttiica (convolution, softplus.tnhh,
_native whwtch_nnrmmmmmooeteeaa .norm.legii........ functional), (convolution. wottioo...otssssooooooo hanpluu.tanh.mul....
native batch norm legit. functional), add]
38
0.9515
[oonvtransposckd, avgpool, clamp. softmax, multiplyl
(convtransposcdd), (avgpool), (clamp, soltmax, multiply!!
[convolution, convolution.
(convolution.avy wow.aaaag-.pp..eaeee. p..aamp softmax, [soffmax, mul)] [gemm, sub.mmol_rhhu)
0.9657
[matmull subernet, multiply, relu] [convrannsossc3d, max, max, sum]
[matnull (subtract, multiply, relui] [constranspose, (max, max, sum)]
78
0.9701
[convolution. convolution, (convolution.mas pool3d with indices). max pool3d_wwil indices, sum] (convolution, convolution, convolution, cconvolution,uu..cok
5
0.9909
(convtransposc2d. subtract, tanh]
3
(convtranspose, (subtract, tanh))
Table 17. Fusion Decisions of Level 2 Problems where I.M-Gnneated code was slower than PyTorch Eager. Here we investigate the programs generated hy DeepSeck-R1 I on Level 2 KernelBenchh Problems, specifically those that are slower than Py Torch Fager (sorted by lower speedup). We list the sequence of operators the program, and fusion decisions made by R1, comparing that with torch. rrch. compi impile le generated kernels.
46

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
M. Alternate Input Representations
explore alternative input representations the KernelBench problems to the Language Model other than a Py Torch reference.
For the PyTorch example as shown in Section C.1 that describes the simple operation a+b used in our baseline prompt.
M.1. Natural language representation
Natural Language might be easiest to start for developer from chat interface, but could suffer from ambiguity, especially for complex operations. We provide specification in natural language that is verbose possible.
This program should demonstrate clement-wise tensor addition. This prorgam takes two input tensors the shape and performs element-wiat addition between them. Each input tensor has batch size of and contains 128 features.
The model has learnable parameters and simply acts as direct addition operation between its inputs. Both input tensors expected to be the GFU (CUDA device).
Operation: klement-wise addition
Input shapes: (1, 128) for both inputs
Output shape: (1, 128)
Parameters: None
10 Device: CUDA (GPU)
11 Purpose: Simple demonstration of tensor addition
M.2. Directed-acyclic graph (DAG) representation
DAG representation explicitly layouts all transforms and operators in the program, which might be more beneficial for the model to conduct kernel fusion (fusing nodes in the graph).
ONNX Graph Representation'
Here is graph representation of the PyTorch program.
Inputs:
I'input', "b'l
Outputs: I' ['add']
Nodes:
Op Type: Add
Inputs: ['input', [b']
Outputs: ['add']
10
Attributes:
12
13
To be specific, the input we will provide will be:
14
tensor of shape (1, 128) of FP32
15
tensor shape (1, 128) of FP32
Torch FX Representation
Here is a DAG representation of the program in tabular form
opcode
name target
args kwargs
placeholder placeholder
()
[}
b
b
()
[}
functiin add <built-in function add>
(a, b)
(}
output
cutput output
(add,)
IONNX (Open Neural Network Exchange) is standard representation neural network programs, which could exported from
PyTorch, https: : /onnx. / /////////// ftnx//ittr
PoyTorch FX is toolkit that help extract symbolic trace and intermediate representations of PyTorch Programs, www.h:ttpt.tt..... http:::/pttttorreeeeeeo
org/docs/stabl attoocsss/sbll...... e/ fx, html
47

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
10
To be specific, the input we will provide will be:
11 a tensor of shape (1, 128) of EP32
12
tensor shape (1, 128) of FP32
M.3. Result Testing Representative Problem
We the same format for the baseline prompt in C.I, instead of updating the task from PyTorch PyTorch to CUDA, Alterative Format (Natural Language, DAG) representation to PyTorch + CUDA. We replace the pair of <PyTorch, CUDA> one-shot example with <Add example in desired Table 18 shows the result when testing this on representative problem in KernelBench, specifically Level 2 19 ConvTranspose2d GELU GroupNorm We chose this example because it was correct when using the PyTorch representation, and represents a semantically common pattern in Level 2. We hope test on this example give us insights on the strengths and weaknesses of various representations.
iinres,etttion,, CUDA>,
Input Representation
Execution Result
Analysis
PyTorch
Correct
Implement Operator in CUDA, No Fusion Wrong initialization, Did not use weight parameters- Implement Operator in CUDA, Wrong Normalization Attempted Fusion, incorrect group convolution
Natural Language
Compilation Error Logical Error
DAG (ONNX)
Output Value Mismatch
DAG (torch FX)
Output Value Mismatch
Table 18. Degraded performance with non-Py Torch input specifications example KernelBench problem. Evaluation of Level 2 problem alternative input representations, best of sampling.
non-PyTorch www.rrrassppsss..coomm Grrrurrroooommm.. using OpenAl GPT-4o and ol (medium reasoning effort), across various
As we observed, when replacing the PyTorch representation with other forms of specifications, the model suffers from ambiguity in its specifications. With natural language especially, despite giving it a verbose description, the model gets confused and creates program that has wrong initialization or neglects to use weight parameters altogether. DAG representations, especially torch FX, helped model to discover fusion; but going from graph to kernel add additional challenge to implement functionally correct program. DAG representations might be beneficial complement to the reference program in PyTorch, though they are not sufficient replacement.
We chose the PyTorch code reference precise, unambiguous, and verifiable input for KernelBench, which is crucial for rigorous benchmarking. Using PyTorch often mirrors the experience of Al researchers (as noted in Section 1, 3) who start with Py PyTorch and then optimize (e.g., popular Attentio (Kim, 2024), Mamba (Dao & Gu, 2024b) repositories follow the workflow of PyTorch code with some inline CUDA kernels). The PyTorch reference is also an unambiguous source for verifying generated code for correctness. It is also worth noting that LLMs are fine-tuned on sequence data, such as code, and PyTorch a natural choice for benchmarking and setup as a code translation task. Future work could consider alternative representations.
48

---

KernelBench:
Can LLMs Write Efficient GPU Kernels?
N. Performance Degradation Analysis
Table 19 presents the best speedup achieved by any evaluated model for each problem instance, sorted first by Level and then by ascending speedup (defined as generated code wall clock time over baseline torch reference time). We note the following observations and include some examples of LLM-generated code:
1. Many foundational operations in level 1, particularly matrix multiplications (matmuls) and convolutions (convs), frequently result in significant slowdowns compared the highly-optimized proprietary libraries like cuDNN used by PyTorch's baseline that use more advanced hardware features. For instance, even the best model achieved only 0.0105 speedup for 69_conv _tt_oov_rranspsseed__. www.:rrnappsddd2..ssmeeiiiiiiiip tensor
and 0. 0.1011 for
tensor matrix multiplication.
2. While models sometimes correctly identify opportunities for fusion in Level 2 tasks, the resulting kernel performance is often poor. This suggests that even if the fusion strategy is conceptually sound, the LLM's implementation of the core fused components (like the matmul or convolution within the sequence) is inefficient, negating any potential benefits from reduced memory access; examples like DeGGmemmmeLel Revide (0.1849 speedup) and 59 39_MMa Swish Scaling (0.1996 speedup) illustrate this issue where complex fusions are much slower than the baseline both cases successfully fuse three operators into a single kernel but they're 're disadvantaged by the matmul implementation.
3. We observe patterns suggesting that LLMs might generate a custom CUDA kernel for a relatively minor operation within a larger sequence. While potentially optimizing that single step locally, this can disrupt Py Torch's native fusion capabilities (like TorchInductor), preventing more holistic, graph-level optimizations across the sequence that could have yielded greater overall performance gains. One example we found is 4 4_LeNet5 in level 3. The model generated custom implementation for ReLU that resulted in 0.6451 speedup.
The generated code for the example problems listed above is included below:
N.1. Code Examples
Note all speedups are against Py Eager.
Slow core operator: 0.0105 speedup for 69_conv_transposed 2D aaymmetriccnnpuutaaynnnee
global__ void conv conv_transposee2__eemme
const float input,
const float weight,
float* output,
const int size,
6
const int in ch__hannels,
const int cut_channels,
const int height_in, const int width_in,
10
const int height_out, const int width_out, const int kernel_h, const int kernel_w, const int stride h,
11
12
13
14
15
const int stride_w,
16
const int padding_h, const padding_w,
17
18
const int cut_padding_h, const int cut pat__padding_ w, const int dilation h,
19
20
21
const int dilation_w const int groups
22
23
)
{
24
const int h_uut = blockIdx.y; const int w_out = blockIdx.xx Tdx.x;
25
26
const int batch_ids blockIdx.2 out out_channe11;
const int cut_ block zxx z out cut_channelss
49

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
28
29
if (hu height_out II w_out width_out || batch_idx batch_size)
30
return;
31
32
float
0.0f;
const int in_ch Leh__pe__grouee in in__hannels / groups;
34
const int out_ch_per_group out_channels / groups;
35
const int group = out_ch / out_ch__err_goupp)
36
for (int kh 0;
kernel_h; kh++)
for (int kw 0; kw < kernel_w; kw++) {
39
const int h_in - (huout + padding_h - kh dilation_h) / stride_hi
40
const int w (w_out + padding_x kw dilation_w) / stride_Mi
41
if (h 0 a2 h in height in as x wiin >= w iin width in) for (int in_ch ch = group * in ch per_ group; in_ch < (group + 1) in_ch_per_grrup; in_ch++)
43
44
45
46
const float input_val input batch_idx in_channels * height_in width_in in_ch height_in width_in h width wiin
48
49
50
52
1;
53
54
const float weight_val weight
55
in_ch cut_ch_per_grouu kernel_h kernel_w
56
(out_ ch 8 out_ ch per groug] kernel kernei_ kh kernel_w kw
57
58
59
1;
60
sum += input_val weight_val;
62
63
64
66
67
output
68
baxch_idx out_channels height_out width_out
69
out_ch height_out width_out
70
hout width_out w out
71
72
] = sum;
Slow operator: 0.1011 speedup for 11.4D.tensor.matrii... muliiiiimuttiiiiitioonn
global void tensor matmul kernel (const float const float float C, int b, int < i, j, int 1, int k)
int idx_b blockIdx.x;
int idx_i = blockIdx.y;
int idx_j blockIdx.2;
int idx_ threadIdx.
if (idx_b < b & idx_i
55 idx_i < & idx k) {
float sum 0.0f; for (int idx sum += Alidx a j * idx idx i idx 11 Blidx 4 + idx kl;
idx 1 ++idx 1]
10
11
12
b
idx
idx idx_k] sum;
13
14
50

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Good fusion pattern but bad implementation: 0.1849 speedup for 63_Gemm.. Ceeemm_ReUU_.. ReLU_Divide
global_ veid linear_res mewearrrelu_e kernel kiekrreel
const flcat* rrsstrict X
const float* restrict_ W,
4
const float* restrict_ D, float* restrict_ out, int batch_size,
int in_features, int out_features, float divisor
10
int idx blockIdx.x blockDim.x threadIdx.x;
12
int total batch_sizo out_features;
13
if (idx < tctal) {
14
int i idx out_fcatures;
15
int j idx $ out_features;
16
float val 0.0f;
17
for (int k 0; k < in in_ k++) {
18
val += xli in teatures kl w wlj in features + kl;
19
20
val += blil; val val 0.f ? val 0.f;
22
val /- divisor; out[idx] val;
23
24
Good fusion pattern but bad implementation: 0.1996 speedup for 59 Matmul Swish Scaling
global_ void lused__inear_swll swish__call__erneel
const flcat input, const fleat* weight, const float+ bias, floats cutput, int B, int IN, int OUT, float scale) {
4
int tid blockIdx.x blockDim.x threadIdx. OUT) return;
(tid
int row tid / OUT; int col tid OUT;
float sum 0.0f; 10
12
for (int 1=0; i<IN; ++i) {
13
sum += input [row=IN IN + i] weight(colII] il
14
15
sum bias[col]s
16
17
float swish - sum (1.0f (1.0f expf (-
18
output [tid] swish scale;
19
Minor custom kernel impacting framework's inherent fusion ability: 0.6451 speedup for 4 LeNet5
relu_source H
#inolude storch/extensee #inalude i h>
global veid relu kernel (const float* in, floatx out, int size) { int idx blockIdx.x blockDim.x threadIdx.x;
if (idx size) {
float val in [idx]; out [idx] = val > 0.0f val : 0.0f;
10
11
12
51

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
13
torsh:: :ensee relu_cuda (torch: Tenser input)
14
auto size input. numel ()
15
auto out - : (input);
16
17
const int block_side size 256; const int grid_siza (size + block_size 1) block_sized
18
19
20
relu celuxkonel<<<grrdd.tt.. block_size>>>
21
imput. dtta out. da..aaaprrrflaat> (int)
22
24
};
25
26
return out;
28
29
30
relu relu souucce I"torch:: relu_cuda (torch: Tensor input) ;"
Compile the inline custom ReLU operator
33
relu _ll__moulle__ee = lcad_inline(
34
name="custco_rellu" cpp_sources=relu_.rr source,
35
ouda www.sscserreuussrrrrrre
37
(ratiins=="relu_uuee verbose-False
38
39
40
class ModelNew (nn :
42
def (self, num num_classes):
43
44
Optimized LeNet-5 architecture with custom CUDA-based ReLU. ***
45
46
super (ModelNew, .elf)._nnite init_
47
48
Convelutional layers
49
self.convl nn.Conv2d (in_chmneeel,, out_channels-6, kerrel_size=5, stride=1) self.conv2 pn. Conv2d (in_cnneei_ce out_chamels==1, sernei_size-t stride=1)
50
51
52
Fully connected layers
53
self.fo nn. Linear fen_fffatures1165555, 1ettures--1)) self.fet fc2 nn. Linear (1n_features-120, out_features-84) self.fot nn. Linear (in_features-64, out_featuree-mmm_clssss
54
55
56
57
Reference to custom ReLU extension self. walf.custem_reett
58
relu_ccuu
60
def (srwf x)
61
convl + custom relu max pool self.convl (x)
62
63
x F.max_pool2d ternel_size-1 stride-2)
65
66
# ccnv2 + custom relu max pool
67
self.
x self.custom relu (x)
69
X m_axx_ool22dx, kernel_size-2, stride-2)
70
71
Flatten
72
x-view(-1, 16+5+5)
74
# fcl + custom relu
75
= self.fc1 (x)
76
self.custom relu (x)
52

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
78
fc2 custom relu
79
X self. fc2 (x)
80
X self. celf.custon_reluu (x)
81
82
final fc
x self.fc3 (x)
84
return x
Table 19: Best achieved speedup per problem instance across evaluated models, categorized by difficulty level and sorted by ascending speedup
Level
Speedup
Problem
Model
0.0105
claude-3.5-sonee
1
0.0316
claude-3.5-sonnet
1
0.0442
35 3r_GrrupNorm.
openai-ol
1
0.0597
68_conv coonv_transposed.2. 3D
3__ioddd___uuuuurreeeeee www.ssuuureeeeeeeeeeeeeee.eeeeee igquareiiputtassym...
deepseek-rl
1
0.0734
40 LayerNorm
llama-3.14055
1 0.1011
11
dcepseek-V3
1
0.1151
claude-3.5-sonnet
0.1196
16 hhatmul.with..tanssooeee
claude-3.5-sonnet
1 0.1198
18 Matmul Matmul.wihh.rasspsodd....
openai-ol
1 0.1332
93 maskee.ccussum cumsum
openai-ol
1
0.1336
77 conv transposed 3D square input square k...
deepseek-rl
1
0.1553
Matmul with large K dimension
claude-3.5-sonnet
1 0.1576
Square matrix multiplication
claude-3.5-somnet
1
0. 0.94
2 Standard matrix multiplication
claude-3.5-sonnet
1
0.1681
97 CosineSimilarityloss
deepseek-V3
1 0.1871
80 co.coov.. standard 2D 3Duuque input asymmetric..
deepseek-rl
1
0.1893
10.3D DeD.teesor... cannsormmatrixxmmm multiplication
claude-3.5-sonnet
0.1895
17 7T_Matuul_wwhh wtw.::ishhtassssed....
deepseek-rl
1
0.1979
13 Mattmul..o..symmettii matrices
claude-3.5-sonnet
0.2083
S_Marmul.withiimegular.sshh
claude-3.5-sonnet
0.2097
02_conv_standard_200 http:/dddde.....uuuuueemmnaaaaoe
deepseek-rl
0.2164
66Lconv_standard_30 http:/ddwd.....aaaaaeee....p..ne
deepseek-rl
1
0.2523
3 3_Batched matrix maaacheedmatti..mmlliiicccc multipiicaiin
deepseek-rl
0.2584
54_conv 34_convsstamdard. _quure input...quare
deepseek-rl
1 0.274
c6_.cov.coo dontwdepthwise depr.eep .2D
deepseek-rl
1 0.3152
24 LogSoftmax
deepseek-rl
1 0.4097
89 cumsum
openai-ol
1
0.4179
7 Matmul with small K dimension
deepseek-ri
1
0.4479
23. 23_Softmax
deepseek-ri
1 0.4807
15 15_Matmul.forJooerrrriiee wwwpouguurrrinnnaaraa.r www.:eerrraanglarmmaaireeeo mattices
deepseek- V3
1 0.5271
81_conv pronn_transposed.113 2D aasampptsedddaammmmmtterree.... btspymmmtric_ippttsquuepee
deepseek-rl
1
0.5497
87 conv pointwise 2D
openai-ol
1
0.5529
14 Matmul for upper triangular matrices
claude-3.5-sonnet
0.5901
Tall skinny matrix multiplication
claude-3.5-son
53

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Table 19: Best achieved speedup per problem instance across evaluated
models, categorized by difficulty level and sorted by ascending speedup (continued)
Level
Speedup
Problem
Model
0.6027
47 17_Sum Sau.rreduciioo...vrra.dddmsssss
deepseek-rl
1
0.6105
19 ReLU
deepseek- V3
1
0.6908
5 Matrix scalar multiplication
deepseek-rl
1
0.6915
48 Mean reduction over a dimension
claude-3.5-somnet
1 0.6953
44 Average Pooling 1D
deepseek-rl
1
0.7117
27 SELU
claude-3.5-sonnet
1 0.7125
26 GELU
deepseek-rl
1 0.7273
22. Tanh
deepseek-rl
1
0.7349
31_ELU
claude-3.5-sonnet
1 0.7375
28 HardSigmoii
deepscek-rl
1
0.7756
20_LeakyReI.U
deepseek-rl
0.7809
deepseek-rl
0.8112
29 9oSSotplus
deepseek-rl
1 0.8117
34 34_InstanceNorm
gpt-4o
1 0.83
39 19_22N.
deepseek-rl
1 0.8622
67_conv.standard ID
deepseek-rl
1 0.8702
21 Sigmoid
deepseek-rl
1 0.872
46 Average Pooling 3D
dcepscek-rl
1
0.8808
32 HardTanh
Blaude-3.5-somer
1 0.8942
85 conv depthwise 2D asymmetric input asymme...
deepseek-rl
1 0.9352
94_MSEL
deepseek- V3
1
0.9439
96 9u_euLsrLsss
deepseek-rl
1
0.9543
4 Matrix vector multiplication.
deepseek-rl
1 0.9671
38 L1Norm
gpt-4o
1 0.9737
90 cumprod
deepseek-rl
1
0.9842
52 Argmin over dimension
deepseek-rl
1
0.9859
45 EnteragePPoolingg 2D
openai-ol
1
1.0022
65 conv transposed 2D square asymm...
openai-ol
1,0084
37 27_FobeniusNorm.
gpt-4o
1 1.0115
98 LLDivLoss
deepseek-V3
1 1.0131
25_Swish
openai-ol
1
1.0259
53
deepseek-rl
1
1.0308
MMax.redutiooo.ovr..aiimmmii
gpt-4o
1 1.0697
100 100_ HingeLssssss
llama-3.1-700
1 1.093
51 Arggax_-ooeee overra.vee...ooo wttp:iwwoveeeradddddee aimmennsoo
openai-ol
1 1.116
42 Max Laox_Pooling22. 2D
openai-ol
1
1.1386
41 Max Pooling ID
openai-ol
1
1.1417
76 conv standard ID dilated strided
deepseck-rl
54

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Table 19: Best achieved speedup per problem instance across evaluated
models, categorized by difficulty level and sorted by ascending speedup (continued)
Level
Speedup
Problem
Model
1
1.1545
43_Max Max Pooling.
deepseek-rl
1
1.185
13_coov_deptie deonhdee 2D square wquare.input anpmmmtsym.trr.......
deepseek-rl
1
1,2706
conv convytransposed... ID
deepseek-rl
1
1.4855
79 conv transposed ID asymmetric input squar...
deepseek-rl
1
1.5
30 Softsign
deepseek-rl
1
1,5319
33 BatchNorm
deepseek-rl
1
1.7083
36 RMSNorm
openai-ol
1 1.9202
95 Lross
claude-3.5-sonnet
1
1.9637
99 Tr_TiipletMaggsnLooss
deepseek-rl
1
4.1341
88 S8.MinGPTNewGelu
deepscek-rl
1
13.2243
12_Mattullywithhdiaaaee
claude-3.5-sonnet
2
1.4849
63 De_.Gemm.Rel.UDiiie
openai-ol
2
0. 1.9996
59
deepseek-rl
2
0.3008
33 Geemm_Scale_aattieee
openai-ol
2
0.3564
94_Gemm Wdd.GemmmBaaAdd...cc hww.aiissdddaadrraahhhhhhha BraittahhMMish.Grruourr
openai-ol
2
0.4642
66 Matmul Matmul MaanoouttMManneSfeeeeaee
openai-ol
2 0.4912
44 ConC Transpoodd Multiply GlobalAvgPoll Glob...
openai-ol
2 0.5501
27 Conv3d HardSwish ReLU Softmax Mean
dcepseek-rl
2
0.5672
67 Conv2d GELU GlobalAvgPool
openai-ol
2
0.6156
80 Gemm Max Subtract GELU
deepseek-rl
2
0.6777
4. Mish
openai-ol
2
0.6922
37 MaatmullSiishhSSSmnrruuop..
deepseek-rl
2
0.7181
85 GrrupNormmSSeeeeee MaxPool Clamp
openai-ol
2
0.7433
69 Conv2d HardSwish ReLU
deepseek-rl
2
0.7451
35 Conv2d Subtract HardSwish MaxPool Mish
openai-o1
2
0.7701
25 Conv2d Min Tanh Tanh
deepseek-rl
2
0.7712
21 C.Cov2d Add. Scale Sigmoid GroupNorm
openai-o1
2
0.7859
36_Conv TonnTrranspoee2d... Min Sin.SSum.GLUU
deepseek-rl
2
0.7915
54_Conv2d
openai-ol
2
0.8231
21_Cony2d.Divid_LLaaay
deepseek-rl
2
0.8377
23.Comv3d.GGooupNorm....
openai-ol
2
0.8399
32.Conv2d.Scaling.M.
deepseek-rl
2
0.8411
68.
M6.Mattul.inn.Sburaat
deepseek-rl
2
0.8595
64. 14..Gemm.LLgSune
deepseek-rl
2
0.8615
deepseek-rl
2 0.8622
86 Matmul Divide.GELU GELU
deepseek-rl
2
0.8715
12 Gemm Multiply LeakyReLU
deepseck-rl
2
0.8801
61 Conv Transpoosood ReLU GroupNorm
openai-ol
55

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Table 19: Best achieved speedup per problem instance across evaluated
models, categorized by difficulty level and sorted by ascending speedup (continued)
Level
Speedup
Problem
Model
2
0.9006
46 SoonvddSSubrrace Taniracct.Tanh.... Saa.rraat.Tan..Suuaaaaaaa Auvrraattt
openai-ol
2
0.9008
22 Matmul ResidualAdd Cww.dCllmpppoie Logpu...sS
gpt-4o
2
0.9009
97 Matmul BatchNorm BiasAdd Divide
gpt-4o
2
0.9062
47 Conv3d Mish Tanh
deepseek-rl
2
0.9092
82 Conv2d Tanh Scaling BiasAdd Max
deepseek-rl
2
0.9229
87 Conv2d Subtract Subtract Mish
openai-o1
2
0.9236
3. CoCCC ConnTranspose3d Sum LayerNorm AvgPool GELU
openai-ol
2
0.938
52 2d_.Con22d.Aivvvaa Attivation
openai-ol
2
0.9657
9 M.Maaull MMaumullSubrrreetMMlllllyy ReLU
deepseek-rl
2 0.9701
78
deepscek-rl
2
0.9909
5_CowwTraasppose2d.SSbbraat
openai-ol
2
0.9922
openai-ol
2
1.0014
31 3J_Comy2LLMinn.AdddMMlll
deepseek-rl
2
1.008
B5_Com/Transpooe3dLMeee Sett Tanha_aaanaaaaa
claude-3.5-sonnet
2
1.0114
24 44_ConvYdd_Min.Softtnn
openai-ol
2
1.0118
72. P2.Cor/TTapppssSLBBacccnoonmm
BatchNorm Aaw.:oormmAAgPPooAAAyvgPP
openai-ol
2
1.025
41 Gemm BatchNorm GELU GroupNorm Mcan ReLU
openai-ol
2
1.0511
43 Conv3d Max LogSumExp ReLU
dcepseek-rl
2
1.1045
16 Conv Transpose2d Mish Add Hardtanh Scaling
claude-3.5-sonnel
2
1.11
53 Gemm Scaling Hardtanh GELU
deepseek-rl
2
1.1227
1w0,Coon/Transoos3ddClaayy Min
deepseek-ri
2
1.1327
deepseek-rl
2
1.1331
93 J3..oo//Transpose244 Add Min GELU Multiply
gpt-4o
2
1.1376
96 Conv Transpoanpo Multiply Max GlobalAvgPool.
deepseek- V3
2
1.1503
95 Matmul Add Swish Tanh GELU Hardtanh
deepseek-rl
2
1.1508
10 2rassoaasp MaxPool Hardtanh Mean Tanh
deepseek-rl
2
1.1792
50 Con'Trransposedd Scaling Avd Sc...
BiasAdd
deepseek-rl
2
1.1838
83. 23.CCnv3d.
Swov3dd_GrouNNomm MMin Min. CannCllmppDDae Dropout
openai-ol
2
1,1852
60. Dr.Coon/rannssee33d http://ossd.SSSss..hSn. GroupNom GrrooopNoo..aaddSwsstt HardSwish
deepseek-rl
2
1.1976
openai-ol
2
1.2267
2.
A.CowTTuunpoo222BBBB.d..CCu.noo Clamp......
openai-ol
2
1.2419
openai-ol
2
1.2444
19_ConvTTanspponoo
deepseck-rl
2
1,2534
49
www..Taansoseeed.SSooo.oo. SoftttppsedSSSttmaax...m SwftSao Sigmood
deepseek-rl
2
1.2964
30 30.Gemm. GreepoomrouN Hardtanh
deepseek-rl
2
1.3112
55 Matmul MaxPool_Suuie Scale
deepseek-rl
2
1.32
81 Gemm Swish Divide. Divide.Clamp Tanh Clamp
deepseck-rl
2
1,3333
Conv3d Divide Max GlobalAvgPool BiasAdd Su...
gpt-4o
56

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Table 19: Best achieved speedup per problem instance across evaluated
models, categorized by difficulty level and sorted by ascending speedup (continued)
Level
Speedup
Problem
Model
2
1.3736
79
openai-ol
2
1.3981
62 Matmul Www.GGoooooooomaaaaaaaaeeeee
deepseek-rl
2
1.4289
74 Conv C
Muultiply Leaky
openai-ol
2
1.4467
98 Matmul AvgPool GELU Scale Max
deepseek-rl
2
1.4604
39 Gemm Scale BatchNorm
deepseek-rl
2
1,4662
42 Con'Transpose22 GvobalAvgPool BiasAdd LogSu...
deepseek-rl
2
1,5042
29 Matmul _Mih Mish
deepseek-rl
2
1,5456
17
Divide
deepseek-rl
2
1,5745
26_ CooCCon//raassooe3LAAdd Hadrssss HardSwish
deepseek-rl
2
1.6439
88 S.c.rmm.cwuNooom..vvvvvvttttysss
deepseek-rl
2
1.7277
deepseek-rl
2
1.7476
51 Let.emmSourrccccooooyyooooooools
deepseek-rl
2
1.8268
18 X.NNeun SSnn  aa  yyeeeellllhhhp
deepseek-rl
2
1.8839
48 S.w.oaaaSSciigg.TTmmmnalllillaaa
openai-ol
2
2.0192
claude-3.5-somnet
2
2.3173
7 3..Cov3d.RRLUUUIIIC. ReU..ReLLLLLLRULUUUUUUU GELUU GEL-SSiooid BiasAdd
Sigmoid.BiasAddd
claude-3.5-sonnet
2
2.3347
40 Matmul Scallig ResidualAdd
openai-ol
2
2.4595
20 Cony Transpoood Sum ResidualAdd Multiply R...
dcepscek-rl
2
2.6044
14 Gemm Divide Sum Scaling
claude-3.5-somee
3
0.0107
33 VanillaRNN
claude-3.5-somnet
3
0.18
43 MinGPPCCasssllAtttttinn
openai-ol
3
0.3408
34. VanillaRNNHidden
openai-ol
3
0.5782
MMPP
deepseek-rl
3
0.6927
24 EfficiemNNet32
openai-ol
3
0.7214
22 EfficientNetB0
openai-ol
3
0.7382
9 ResNer18
openai-ol
3
0.7455
20 MobileNerV2
openai-ol
3
0.767
23 EfficientNetBI
deepseek-V3
3
0.7815
10. ResNet101
openai-ol
3
0.8083
46 NetVlaaWWihhhhseCsssees
deepseek-V3
3
0.8194
13 DeDnscNNeu////Traiiionaavvv
claude-3.5-sonnet
3 0.8249
22_CCovouttomiisiiinTaaaffooo
openai-ol
3
0.8284
7 7.GooglaNeelnccptioon\.
openai-ol
3
0.855
15 DenseNet121
gpt-40
3 0.8587
18 SqueezeNet
gpt-4o
3
0.8591
16 DenseNet201
gpt-4o
3
0.8699
47 NeWVlaNooGhosShusters
1lama-3.1-405b
3
0.8743
28 VisionTransformer
gpt-4o
57

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
Table 19: Best achieved speedup per problem instance across evaluated
models, categorized by difficulty level and sorted by ascending speedup (continued)
Level
Speedup
Problem
Model
3
0.8846
19 MobileNetVi
deepseek-V3 V3
3
0.9343
27 RegNet
openai-ol
3
0.9608
26 ShuffleNet
gpt-4o
3
0.9811
AlexNet
openai-ol
3
0.9858
25 ShuffleNetUnit
deepseek-V3
3
1.0048
14 DenseNNe122/DenseBBlec
openai-ol
3
1.0089
6 GoogleeNetncceiionMoooulee
openai-ol
3
1.012
ResNetBasicBlock
openai-ol
3
1.025
29 SWWnMLL
llama-3.1-4056
3
1.027
11_VGG16
dcepseek-V3
3
1.0442
12_VGG19
deepseek- V3
3
1.08
36J.TSMHH
openai-ol
3
1.0922
3 D.De
deepseek-rl
3
1.2704
44 M4.MiniGPTBlock
deepseek-V3
3 1.405
4_LeNe55
claude-3.5-sonnet
3
1.9376
50. Re.eSeelAAttnnnonnn
openai-ol
O. A Alternative Library: Triton
Alternative GPU programming tools to CUDA such as CUTLASS (NVIDIA, 2017a), Triton (Tillet et al., 2019), Thun- derKittens (Spector et al., 2024) have been developed to make GPU programming easier by exposing a higher level of abstraction. We agree that this is exciting direction to improve model performance and had noted this direction for future work in Section 6.2. We reiterate that the goal of this work is to propose new benchmark framework and to thoroughly evaluate the baselines, rather than to solve the full kernel generation problem.
0.1. Triton Task Specification
We extend KernelBench with Triton evaluation backend and Triton task specification. Instead of <PyTorch, PyTorch + CUDA>, we now define the task as <PyTorch, PyTorch Triton>.
Similar to the PyTorc one-shot example used in C.I that describes the simple operation a+b, we provide the model with an inline Just-In- Lust-en-Time compilation (JIT) kernel example in Triton.
import torch
import terch.nn as import torch.nn. functiona as import triton
import triton. language as tl
7
Striton. jit
def add_ker
10
xtr,, $ Pointer to first input y_ptr, & Pointer to second input
11
12
pur, # Pointer to output
13
n_elements, Total number of elerents in input/utppt
14
BLOCK_STZE cl.constexpr,
58

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
15
):
16
Each program handles a contiguous block of data of size BLOCK_SI2E block_start - tl. patoprrranmidd (0) BLOCK_SI&E # Create range of offsets [0.. BLCCK SI2E-1]
17
18
19
offsets block_start
tr.anange (0, BLOCK_SIZE)
Mask to ensure we don't go out bounds mask = offsets n_elements
21
22
Load input values
23
(1.1cad(x_pt)
offsets, mask=nask, other=0.00
24
y ty_ppr offsets, mask=mask, other=0.0)
Perform the elementwise addition out x +
26
27
Store the result
28
tt.store (out_ptr offsets, out, mask=mask)
30
31
def triton_add (x: torch. Tensor, torch. Tensor):
32
***
33
This function wraps the Triten kernel call. It:
34
1. Ensures the inputs are contiguous on GPU.
35
2. Calculates the grid (blocks) needed.
36
3. Launches the Triton kernel.
37
***
assert x.is_cuda cuda and y.is_cuda, "Tensors must be on CUDA.
39
x (:contiguous()
40
y.contiguous()
41
42
Prepare cutput tensor out terch. enety_elptlllleee like (x)
44
45
Number of elements in the tensor n_elements n.mumll
46
SIZE 128 Tunable paraneter for block size
48
49
Determine the number of blocks needed grid = lambda meta: ((n_elements meta ["BLOCK_SI2="] 1) // meta "BLOCK_STZE")
50
51
Launch the Triton kernei
53
Mel.kernellgge (x, out, a elements, BLOCK SIZE__LIIEBBLCK SIZE) return out
54
55
56
57
class ModelNew (nn Module)
58
def _intt__ii (self) -> None:
59
super _uper()__nit___. ()
60
def (sclf a, b);
62
Instead of "return a + b", call our Triton-based addition return triton_add b)
63
fast, over:
Py Torch Eager
torch.compile
KernelBench Level
GPT-4o
2%
7%
2%
13%
2%
2%
OpenAl ol
3%
17%
10 %
22%
13%
8%
DeepSeek R1
6%
13%
2%
19%
13%
4%
Llama 1-7BB Inst.
1%
0%
0%
9%
0%
0%
Table 20. BernelBench-Triton Baseline. Similar to Table 1, we present fast, against PyTorch Eager and compile NVIDIA L40S. The LMS target Triton here instead of CUDA.
59

---

KernelBench: Can LLMs Write Efficient GPU Kernels?
O.2. Baseline Evaluation with Triton
We again conduct baseline evaluation on Levels on KernelBench, just like Section 4.1, but with the Triton task and associated evaluation. As shown in Table 20, find that models perform worse when targeting Triton compared to CUDA (Table 1), both in terms of correctness and performance. For example, DeepSeek RI fast; drops from 12%, 36%, 2% to 6%, 13%, 2% for Level 1, 2, and 3 respectively. Analyzing the errors and samples, we found many Triton-related errors, likely due to Triton being a more rare source of training data than CUDA, highlighting challenges for using low-resource domain-specific libraries.
60


Reviews:
[{"review_id": "ZHOh0r1tYp", "reviewer": "Reviewer_AoMw", "summary": "This paper proposes KernelBench, which is a new benchmark for evaluating LLM's performance in writing correct and fast kernels. Specifically, KernelBench gathers three different levels of tasks, including individual operations, sequence of operations, and end-to-end architectures, and introduces a novel fast_p metric to model correctness and efficiency at the same time. KernelBench shows that most frontier models do not perform well in writing kernels, among which the state-of-the-art reasoning models perform the best. KernelBench further shows that leveraging feedback is important for reducing execution errors and discovering faster solutions:", "questions": "", "limitations": "", "rating": 3, "confidence": null, "soundness": null, "presentation": null, "contribution": null, "strengths": "", "weaknesses": "", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "Updated Paper: https://storage.googleapis.com/anonymous-files/kernelbench.pdf\n\nWe thank reviewer AoMw for your review\\! We are glad the reviewer appreciates our \"fast\\_p evaluation metric for KernelBench\" and finds that our \"claims are supported by clear and convincing evidence.\" Below, we address your comments regarding the evaluation metrics, platform dependency, and comparison to other benchmarks.\n\n**Clarity of fast\\_p vs. Separate Correctness and Speedup Metrics**\n\n“Adding two separate metrics for correctness and speedup” is a great suggestion for added clarity. In our revision (see link above), we included an extended table presenting correctness (equivalent to fast\\_{p=0}) and geo-mean speedup as separate metrics (Appendix I, Table 16), providing a disaggregated view for these specific aspects. The geometric mean of speedups only includes the correct generations, as fast but incorrect code is not helpful. \n\nThanks for acknowledging that the “evaluation metric fast\\_p makes sense for KernelBench\\!” We also reiterate that for kernel generation, speedup and correctness are tightly coupled, motivating our choice to design fast\\_p. To provide more context on this, we've also added a section discussing our metric design explorations (Appendix I). We hope the combination of the original fast\\_p and these new separate metrics offers more clarity.\n\n**Platform Dependency and Reproducibility**\n\nWe agree entirely about the platform-specific nature of hardware performance tuning. Thus, it is very important to compare results when controlling for both the input program and the underlying hardware. For instance, our evaluations (as noted in paper Section 4.4 and Appendix G.1) across several hardware platforms, including L40S, A100, and H100 GPUs, revealed reasonable consistency in the kernel generations at Level 1 but more pronounced variation in Level 2\\. Other than this hardware evaluation study, most of our experiments in the paper are done on an Nvidia L40S, and we expect the results to be reproducible on this type of GPU.\n\n**Comparison to Existing Programming Benchmarks**\n\nTo address your comment, we've added a new experimental section (see Appendix J) where we compare model performance on KernelBench (KB) with LiveCodeBench (LCB). (We chose LCB as HumanEval performance is quite saturated for current models). We present the relative rankings of models across these benchmarks. \n\nOur results show, perhaps unsurprisingly, that models performing well on general coding benchmarks tend to also perform better on KernelBench, but variability in rankings (e.g. o1 ranks 1st in LCB and 2nd in KB Level 1, and R1 ranks 2nd in LCB but 1st in KB Level1) across different levels of KernelBench suggests that additional skills are required for high performance in kernel-specific tasks – intuitively, this aligns with the major differences between GPU programming and standard programming problems found in popular coding benchmarks. We would also like to highlight that KernelBench is not merely another code generation benchmark; it adds the critical dimensions of performance optimization and hardware awareness, testing a model's ability to generate not only correct code, but also efficient code, which presents distinct challenges. \n\n**Relation to Broader Literature of Leveraging Feedback**\n\nWe definitely see the connection with existing works on leveraging feedback too (we also added citations of works listed by you here)\\! We believe KernelBench takes this concept into a particularly challenging and impactful domain. Optimizing hardware kernels (notoriously difficult even for human experts) offers tangible real-world benefits (cost and energy savings for AI\\!), making it a high-stakes environment given the ubiquity and importance of AI systems today. \n\nWe intentionally designed KernelBench to facilitate precisely iterative, feedback-driven improvement. By providing rich, actionable feedback signals—clear correctness checks (pass/fail), compilation status, precise runtime measurements, and speedup relative to a baseline—KernelBench creates an environment where AI systems can directly learn from their attempts and refine their solutions. The goal of our baseline results using feedback is to thoroughly characterize the degree to which we can solve KernelBench. We find that despite using these techniques, the best model gets fast\\_p=1 of only 18% on level 3, showing there’s a lot more progress to be made on this benchmark. In this sense, we see KernelBench as a stepping stone for pushing forward research in automated kernel engineering, providing a crucial contribution to the community as a standard evaluation environment.\n\nIn light of these clarifications as well as new experimental results and modifications to the paper to address your comments, we would really appreciate it if you would re-examine our paper and consider raising your score.", "claims_and_evidence": "Most claims made in the submission are supported by clear and convincing evidence.", "methods_and_evaluation_criteria": "The evaluation metric fast_p makes sense for KernelBench. However, apart from fast_p which takes both correctness and speedup into consideration, I would recommend adding two separate metrics for correctness and speedup separately for clearer demonstration.\n\nFor the evaluation approach, KernelBench does not handle the issue of cross-platform variation. While the paper claims that it \"does not provide ground truth kernels for the tasks since we imagine users benchmarking on a variety of hardware platforms (including new platforms)\", it's possible that, in terms of speedup, some model performs the best on one platform while the other model performs the best on the other platform -- it will make the evaluation results (even the model rankings) super hard to reproduce.", "theoretical_claims": "There's no theoretical claim in this paper.", "experimental_designs_or_analyses": "There is one weakness regarding the analysis:\n- There is a lack of comparison to existing programming benchmarks like HumanEval/MBPP and LiveCodeBench. While it is reasonable that some models are expected to have various ranks across different benchmarks as KernelBench has a different focus, the overall ranking should largely align well with the mainstream benchmarks.", "relation_to_broader_scientific_literature": "This paper shows that leveraging feedback is important for reducing execution errors and discovering faster solutions, which aligns well with many existing works that focus on more general code generation tasks [1,2].\n\n[1] Chen, Xinyun, et al. \"Teaching large language models to self-debug.\" arXiv preprint arXiv:2304.05128 (2023).\n\n[2] Xia, Chunqiu Steven, and Lingming Zhang. \"Keep the Conversation Going: Fixing 162 out of 337 bugs for $0.42 each using ChatGPT.\" arXiv preprint arXiv:2304.00385 (2023).", "essential_references_not_discussed": "N/A", "other_strengths_and_weaknesses": "N/A", "other_comments_or_suggestions": "N/A", "questions_for_authors": "N/A", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}, {"review_id": "KlvrLpf5zF", "reviewer": "Reviewer_ufH2", "summary": "This paper introduces KernelBench, a benchmarking framework designed specifically to evaluate the correctness and performance of GPU CUDA kernels generated by large language models (LLMs). KernelBench compiles a representative set of PyTorch code snippets, categorizing them into three distinct complexity levels based on their granularity. By systematically assessing the correctness and performance gains achieved by various LLMs when converting these PyTorch snippets into CUDA kernels, KernelBench provides a comprehensive evaluation of the kernel generation capabilities across multiple prominent LLMs and hardwares.", "questions": "", "limitations": "", "rating": 3, "confidence": null, "soundness": null, "presentation": null, "contribution": null, "strengths": "", "weaknesses": "", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "Updated Paper: https://storage.googleapis.com/anonymous-files/kernelbench.pdf\n\nWe sincerely thank you for the detailed and insightful review\\! We are truly encouraged by the positive feedback, particularly that the work is seen as \"well-supported by strong evidence\" and a \"comprehensive benchmark.\" We especially appreciate the recognition of our \"novel metric, fast\\_p,\" as offering a \"more holistic and nuanced assessment,\" and that you found value in our \"wide range of experiments across multiple dimensions.\" We were glad to read that you felt our work \"provides valuable insights into \\[LLMs'\\] capabilities in generating CUDA kernels\", which are valuable for AI and HPC communities as we explore automating kernel optimization. These strengths highlight core goals we aimed for, and we're glad they resonated.\n\n**Analysis of Performance Degradation Cases**  \nPer your suggestion, in addition to our existing error analysis (4.2) and case study of the fastest kernels (Appendix D), we added a new “Case study: Performance Degradation” (Appendix N), which specifically examines instances where generated kernels underperformed compared to the baseline. Here are our findings:\n\n1. LLM implementations of core ops (matmul, conv) underperform highly-optimized proprietary (e.g. cuDNN) kernels in PyTorch.  \n2. LLM correctly identifies fusion patterns, but fused operations (often matmuls) are not efficiently implemented, outweighing benefits from reduced memory access  \n3. LLM blocks better PyTorch native fusion capabilities by generating a custom kernel for a minor task that prevents optimizing across a larger sequence of operations.\n\n**Alternative Input Specification**  \nPer your suggestion, we explored in Appendix M using 1\\) Natural Language 2\\) DAGs of program operators as input specification. On a representative Level 2 problem that model succeeded with PyTorch representations, it failed with compilation and logical issues on natural language representation due to ambiguity on exact behaviors, even when provided with verbose dimension details. DAG representations capture the program execution much better and hint the model to conduct fusions, but going from DAG to kernel directly can lead to logical errors resulting in output mismatch.\n\n**Using Libraries & DSLs**   \nReviewer raises the point of whether generating code using frameworks like Triton/CUTLASS would be helpful. To address your feedback, we extended KernelBench with a Triton task specification and evaluation backend. As shown in Appendix O Table 20, we found models perform worse when using Triton, both in terms of correctness and performance: fast\\_1 for DeepSeek R1 drops from 12%, 36%, 2% to 6%, 13%, 2% across 3 levels respectively. Qualitatively, we found models generate many Triton-related errors, likely due to Triton being a rarer source of training data than CUDA, highlighting potential challenges for using domain-specific libraries.   \nWe reiterate KernelBench’s goal is to propose a new benchmark with thorough baseline evaluation as a first step, rather than solving kernel generations.\n\n**Level-Specific Prompting and Scoring**\n\nWe made a deliberate choice to not explicitly weight fast\\_p by task complexity (e.g., Level 1 vs. 3). We report the levels separately, and harder tasks are expected to yield lower scores naturally, reflecting greater challenges. This approach is common in other coding benchmarks (LiveCodeBench) that include easy/medium/hard problems without score normalization by difficulty. Regarding level-specific prompting, our baseline evaluations intentionally used general prompts, rather than the suggested task-specific ones, to evaluate each model's fundamental ability to independently discover and select optimization strategies across task complexities without explicit steering.\n\n**Architecture-Specific Optimizations**  \nBased on your suggestions, we added experiments for eliciting architecture-specific optimizations: Tensor Cores and asynchronous memory transfers (Appendix G.3) on Ampere GPUs, in addition to existing experiments in Section 5.2. We provided DeepSeek-R1 with examples using wmma and memcpy\\_async instructions, on simple KernelBench matrix multiply problems in FP16 (compatible with Tensor Cores). We observed that the model attempted to apply yet struggled to utilize those advanced instructions. Among the 5/17 correct kernels that use WMMAs, successfully leveraging Tensor Cores did not lead to better performance over PyTorch. No kernels used memcpy\\_async correctly. This highlights that successfully utilizing hardware features remains challenging for models and KernelBench provides a playground for the community to develop methods that address this limitation. \n\nOnce again, we thank you for your valuable time and feedback. We hope that our additional experiments, analysis of results, and discussion addressed your concerns. We hope you find the paper significantly improved and consider reflecting this in your final score.", "claims_and_evidence": "While the majority of the claims in this paper are well-supported by strong evidence, there are areas where the arguments and evidence could be further strengthened:\n\n- This paper claims to propose a benchmark for evaluating the performance of LLMs in generating CUDA kernels. While the evaluation and analysis are indeed quite detailed, the results show numerous instances of performance degradation or no change. Providing a more comprehensive analysis of cases where LLM-generated kernels result in performance degradation would offer a more balanced and realistic portrayal of LLMs' capabilities in this domain.", "methods_and_evaluation_criteria": "**Methods and Potential Issues:**\n\n- This paper introduces KernelBench to evaluate the kernel generation capabilities of LLMs. However, the design of KernelBench primarily focuses on assessing the translation of PyTorch code into CUDA kernels, which essentially evaluates code translation abilities. Including a broader range of code or natural language to CUDA kernel translations would significantly enhance its value.\n- KernelBench appears to be more suited for generating raw CUDA C++ code. In reality, LLMs might have the potential to leverage other general-purpose tools (such as Triton, CUTLASS) for kernel generation. Exploring these possibilities in the paper would make the results more comprehensive.\n\n**Evaluation Criteria and Potential Issues:**\n\n- This paper employs the formula (in Line 213) as the standard for evaluating LLM-generated GPU kernels:\n\n  This formula integrates both correctness and performance dimensions to comprehensively assess the capabilities of LLMs. However, the metric does not account for differences in task complexity. For example, Level 1 tasks (single operations) and Level 3 tasks (full architectures) vary significantly in difficulty, but the $fast_p$ metric does not weight or adjust for these differences.", "theoretical_claims": "The paper did not make any theoretical claims or proofs.", "experimental_designs_or_analyses": "I have reviewed the experimental designs and analyses for their soundness and validity, and I found no major issues or concerns to address.", "relation_to_broader_scientific_literature": "N/A", "essential_references_not_discussed": "N/A", "other_strengths_and_weaknesses": "**Strengths:**\n\n- The paper introduces a comprehensive benchmark, KernelBench, designed to evaluate the ability and potential of Language Models (LMs) in generating efficient GPU CUDA kernels. \n- The authors propose a novel metric, $fast_p$, which combines both correctness and performance to evaluate the quality of generated kernels. This dual-focus approach offers a more holistic and nuanced assessment compared to traditional metrics that rely solely on correctness or performance. The analysis of mainstream large models using this metric provides valuable insights into their capabilities in generating CUDA kernels.\n- The paper conducts a wide range of experiments across multiple dimensions, including prompt content, hardware types, and operator categories.\n\n**Weaknesses:**\n\n- The paper focuses exclusively on evaluating LMs' ability to generate raw CUDA C++ code. However, it does not explore the potential of LMs to leverage high-performance libraries such as CUTLASS or Triton, which could significantly enhance kernel performance. Investigating these tools could reveal the upper bounds of LLM-generated kernel performance and provide a more comprehensive evaluation.\n- While the paper examines the impact of different GPU parameters, it does not delve into architecture-specific optimizations. Modern GPUs, such as those based on Ampere or Hopper architectures, offer unique features like asynchronous memory access and warp specialization. Incorporating experiments that utilize these architecture-specific optimizations could unlock further performance gains and provide a more complete picture of LLM capabilities.\n- The paper does not address the design of different data types, which are closely related to Tensor Core utilization in CUDA. The benchmark also lacks designation about specific Tensor Core code. Given the importance of Tensor Cores for mixed-precision models (e.g., FP16, FP8) commonly used in SOTA LLMs, incorporating Tensor Core-specific designs would significantly enhance the relevance and value of the benchmark.", "other_comments_or_suggestions": "The optimization prompts for different levels of tasks may need to be differentiated. For instance, for Level 1 tasks (operator-level generation), it might be necessary to guide the LLM to generate lower-level CUDA code, essentially rewriting the entire operator. On the other hand, for Level 3 tasks (graph-level generation), it may not be necessary for the LLM to generate CUDA code for all operators. Instead, it might be more effective to guide the LLM to rewrite part of specific PyTorch functions within the network to achieve optimization.", "questions_for_authors": "See Weakness and Comments.", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}, {"review_id": "dmneIPkprf", "reviewer": "Reviewer_qSdx", "summary": "The major contribution of this paper is as follows:\n- This paper introduced a benchmark framework to evaluate how good a modern LLM can write efficient GPU kernels. The core of this benchmark framework consisting of 250 tasks with 3 levels of granularity: single primitive, sequence of ops and the overall model.\n- Using the benchmark, this paper evaluated modern LLMs. The conclusion is that there's still a large room of improvement as modern LLMs face challenges on both correctness ratio and efficiency. Experiments shows repeated or iterative sampling may help, but still yielding unsatisfactory kernel performance.", "questions": "", "limitations": "", "rating": 3, "confidence": null, "soundness": null, "presentation": null, "contribution": null, "strengths": "", "weaknesses": "", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "Updated Paper: https://storage.googleapis.com/anonymous-files/kernelbench.pdf\n\nWe thank you for appreciating KernelBench design and suggesting further improvements. As you noted, automatic GPU code generation is an underexplored area with many interesting research questions; KernelBench facilitates research in this direction as the first benchmark and environment for kernel development, with “simple-yet-effective” task definition and “well-designed evaluation”. In fact, we have already seen enthusiasm from the community, with multiple projects tackling KernelBench through agentic optimization and post-training.\n\n**Providing Model with Hardware Information**  \nWe totally agree that specifying hardware information is important as GPU kernels are inherently platform-dependent. In fact, this was already studied in our original submission. In Section 5.2.2, we provided the model with the exact kind of GPU hardware specifications (see Appendix C.5) that the reviewer described, and found that current models rarely conduct optimization correctly for underlying hardware when provided with such information. \n\n**Clarifying Design Choice to Test Kernel Fusion**   \nIn KernelBench the model has full flexibility to decide what subset of operators in the PyTorch reference to optimize and fuse. We believe this is one of the crucial abilities when a model is given distinct or new architectures in the real-world setting.\n\nKernelBench’s 3-level categorization helps disentangle fusion decisions and kernel generation. Level 1 problems (single operators) only test the model's ability to write optimized kernels; Level 2 and 3 problems are designed to additionally evaluate the model's ability to identify and leverage fusion opportunities; Appendix K provides a detailed task breakdown. \n\n**Fusion Patterns in Model Generated Code**  \nTo answer your questions about **fusion patterns** in generated kernels, we manually inspected the kernels generated by the best performing model, DeepSeek R1, and provided new analysis in Appendix L. We focus on level 2 problems, which are composed with one mainloop (e.g. conv, matmul) and 2-5 epilogue operations (non-linearities, reductions, etc). We observe model-generated code always attempts to generate 1-2 fused kernels per problem. As shown in Figure 19, the fused kernels tend to contain more than half the operators in the program. To explicitly answer your question, only 18% of programs fuse all operators into a single kernel.\n\nRegarding your question on the quality of fusion decisions and whether they cause low performance, we analyzed the generated kernels that were slower than PyTorch Eager (as shown in Table 17\\) and draw two observations: 1\\) main loop operators (eg. Conv) were not fused with epilogue operators 2\\) the model’s attempt to fuse main loop operators (e.g. GEMM \\+ other ops) was not faster than launching highly-optimized cuBLAS kernels. Also refer to **“Analysis of Performance Degradation Cases”** in response to Reviewer ufH2 for a related study.\n\n**Comparison with SOTA Compiler-Based Approach**  \nThank you for bringing up relevant compiler-based approaches (AStitch, Welder, ROLLER, TVM), which we have added and elaborate on in our updated related works.\n\nTo directly address your concern, we compare fusion decisions in model-generated kernels with auto-fusion compilers. Since AStitch, Welder, ROLLER could not be run on KernelBench due to format incompatibilities or outdated support for KernelBench PyTorch 2.5 /CUDA 12.4 (Appendix B), we focus comparison on widely-adopted torch.compile with SoTA performance (better than TVM, see Table 3 in PyTorch 2 \\[1\\], ASPLOS ‘24) and employs an auto-fusion policy over TorchInductor’s define-by-run IR. We show both fusion decisions of R1 and torch.compile in Table 17\\. Torch.compile often creates sophisticated fusion patterns by breaking Convolutions or GroupNorm into smaller multi-pass kernels that compute partial results and statistics in parallel — behavior that R1-kernels rarely exhibit. \n\n**DAG Representation**  \nPer your suggestion, we conduct experiments on using a DAG representation (ONNX, torch.fx graph) “instead” of a PyTorch Reference, which might help highlight fusion opportunities. We explored this in Appendix M.2 and found that DAG representations cause output mismatch issues on problems that succeed with PyTorch representations – see response to reviewer ufH2 for details.\n\nWe appreciate the comprehensive comments and hope that our additional experiments, analysis of results, and discussion addressed your concerns. We hope you find the paper significantly improved and consider reflecting this in your final score.\n\n\\[1\\] PyTorch 2: Faster Machine Learning Through Dynamic Python Bytecode Transformation and Graph Compilation. ASPLOS '24. https://doi.org/10.1145/3620665.3640366", "claims_and_evidence": "The evaluation in this paper well supports the claim that modern LLMs underperform.", "methods_and_evaluation_criteria": "The design of the benchmark successfully cover basic aspects including throughput, correctness while describing the tasks with plain text + pytorch ref code + sample inputs. This design is simple-yet-effective.\n\nIn several aspects, the paper should have explored deeper.\n- One major concern is letting LLM to generate kernel without hardware information. In real world senarios, the hardware specs like number of registers, shared memory, cuda compatibility, etc would critically affect the performance of a kernel.\n- In the task designed by this paper, it's important for the LLM to decide when and how to select a subset of ops to fuse (write kernel) for. E.g. fusing the attention block, but not the entire transformer block or even the entire model. This paper does not show much information on this and it's unclear whether the low performance of LLMs, especially in task 3, are writing low-performance models because of poor CUDA optimizations or simply writing kernels for blocks that should't be fused.", "theoretical_claims": "Benchmark paper, no theoretical claims.", "experimental_designs_or_analyses": "The evaluation for the correctness and evaluation is well designed.\n\nOne concern mentioned above is the lack of evaluation on whether the decision on the subset of ops to be fused is correct.\nMoreover, it would be beneficial to compare with some SoTA auto fusion methods.", "relation_to_broader_scientific_literature": "This benchmark would facilitate the research on automatic GPU code generation, which is underexplored but attracting increasing attention.", "essential_references_not_discussed": "There's a line of research on automatic kernel generation solving the same problem but using compiler-based methods like AStitch (ASPLOS'22), Welder (OSDI'23), ROLLER (OSDI'22) or older foundation works like TVM (OSDI'18)", "other_strengths_and_weaknesses": "Strengths: A benchmark measuring basics of LLM GPU code generation\n\nWeaknesses: See methods and experiment section", "other_comments_or_suggestions": "- It will be good for the paper to provide more details on the designed tasks, especially for level 2&3 for readers to understand the coverage of this tasks.", "questions_for_authors": "- In the benchmark design, how does the LLM make decision on which part of the computation we need to generate kernel for? Do they tend to generate a single kernel for all computation?\n- Will it be more beneficial to use a DAG representation instead of a pytorch ref code representation?", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}]

Author rebuttal:

Meta-review: The authors give a benchmark for LLM-generation of GPU code, which is doubly relevant to the machine learning community both because of the target application and the method. Reviewers are unanimous in recommending acceptance. I agree with the reviewers and would urge the authors to incorporate discussion of the related work that was raised by qSdx.
Decision: Accept (poster)