PaperlayerView exact revision on arXiv

Source-linked AI summary

CudaForge: An Agent Framework with Hardware Feedback for CUDA Kernel Optimization

Zijian Zhang, Rong Wang, Shiyang Li, Yuebo Luo, Mingyi Hong, Caiwen Ding

arXiv:2511.01884v2cs.LGcs.AIcs.CLcs.DC

TL;DR

Existing LLM approaches struggle to produce CUDA kernels that are both correct and performant, while developing highly efficient kernels requires substantial GPU expertise. CudaForge uses a training-free Coder–Judge workflow that iteratively incorporates hardware feedback; on KernelBench, it reports 97.6% correctness and outperforms existing baselines.

  • Problem

    LLMs still struggle to produce CUDA kernels that are correct or performant out of the box, while high-efficiency kernel development requires deep GPU-architecture and parallel-programming expertise.

    Source: §1 · pp. 1–2

  • Method

    CudaForge uses a training-free workflow in which a Coder and Judge iteratively generate, correct, and optimize kernels using GPU specifications and Nsight Compute feedback.

    Source: §1 · p. 3

  • Results

    Across 250 KernelBench tasks, CudaForge outperforms Kevin-32B, OpenAI-o3, and the Agentic Baseline, with 97.6% correctness.

    Source: §1 · pp. 3–4

  • Takeaways & Limitations

    The authors conclude that CudaForge offers a practical, low-cost approach to developing performant kernels across GPU architectures and base models.

    Source: §1 · pp. 3–4, §6 · p. 18

    One comparison uses published Level 1–2 results because the relevant code was not open-sourced.

    Source: §3.3 · p. 11

Abstract

from arXiv · show

Developing efficient CUDA kernels is increasingly critical for AI applications such as large-scale LLM training. However, manual kernel design is both costly and time-consuming, motivating automatic approaches that leverage LLMs for code generation. Existing methods for automatic kernel generation, however, often produce low-efficiency kernels, incur high computational overhead, and fail to generalize across settings. In this work, we propose CudaForge, a training-free multi-agent workflow for CUDA kernel generation and optimization. Our workflow is inspired by the iterative workflow of human experts, which contains steps such as developing initial kernels, testing correctness, analyzing hardware feedback, and iterative improvement. More specifically, CudaForge employs two LLM agents: a Coder and a Judge, that iteratively generate, correct, and optimize CUDA kernels, while integrating hardware feedback such as Nsight Compute (NCU) metrics. In extensive evaluations, we show that CudaForge, by leveraging base models like OpenAI-o3, achieves 97.6\% correctness of generated kernels and an average 1.68$\times$ speedup over PyTorch baselines, substantially surpassing state-of-the-art models including OpenAI-o3 and Kevin on KernelBench.Beyond accuracy and speed, CudaForge demonstrates strong generalization across GPUs (A100, RTX 6000, 4090, 3090) and base models (OpenAI-o3, GPT-5, gpt-oss-120B, Claude-Sonnet-4, QwQ-32B), while maintaining high efficiency. In particular, generating an optimized kernel takes about 26.5 minutes on one RTX6000 and incurs about \$ 0.3 API cost, which is significantly cheaper than existing agentic work that costs 6 H100 hours and \$ 5 API cost per kernel. Our results highlight that multi-agent, training-free workflows can enable cost-effective, generalizable, and high-performance CUDA kernel optimization. Code available at https://github.com/OptimAI-Lab/CudaForge

1 Introduction

Generating efficient CUDA kernels is difficult, and existing LLM approaches face limits in kernel efficiency, cost, and hardware-aware guidance. CudaForge introduces a training-free Coder–Judge workflow using hardware feedback; it reports strong KernelBench results and low-cost operation.

1 Introduction

  • 97.6% correctness and 1.68× average speedup over PyTorch baselines on 250 KernelBench tasks, surpassing Kevin-32B, OpenAI-o3, and the Agentic Baseline.Figure 1 presents the KernelBench comparison.Source: Figure 1 · p. 2, §1 · pp. 3–4
  • Kevin-32B averages 1.10× speedup over KernelBench Levels 1–2, while RL methods require substantial resources and existing agentic work costs about 6 H100 hours and $5 per kernel.The introduction also notes that some models struggle to produce correct or performant kernels out of the box.Source: §1 · pp. 2–3
  • CudaForge separates kernel generation and evaluation between a Coder and Judge, integrating GPU specifications and Nsight Compute metrics to identify bottlenecks and guide targeted optimization.The workflow follows iterative steps modeled on human kernel development.Source: §1 · p. 3
  • The training-free workflow averages 26.5 minutes and about $0.3 API cost per kernel on one RTX 6000, and the authors report robustness across GPU architectures and base models.The paper also reports that increasing iteration rounds can further improve performance.Source: §1 · pp. 3–4

2 The CudaForge Framework for CUDA Kernel Optimization

CudaForge iteratively generates and tests CUDA kernels, then uses runtime and hardware feedback to correct errors or target performance bottlenecks. Its design combines specialized agents with an offline-selected metric set and focused, round-by-round optimization guidance.

2.1 CudaForge Framework

  • The Coder generates a candidate, the Judge returns correction feedback for failing kernels or optimization guidance for correct ones, and the workflow selects the most efficient correct kernel after up to N rounds.Profiling uses NCU metrics together with GPU specifications to identify bottlenecks.Source: §2.1 · pp. 4, 6

2.2 Component Design

  • The Coder receives a one-shot code example and only the latest round context; kernels must compile and match PyTorch outputs within 1e−4, while the Judge switches from correction to optimization after validation.The Judge returns structured feedback for the next generation round.Source: §2.2 · pp. 6–7

2.3 How to Integrate Hardware Feedback

  • At each optimization round, the Judge focuses on 3–4 metrics alongside GPU specifications, identifies a dominant bottleneck, and gives targeted guidance for the Coder to address it.The workflow directs each round toward one critical bottleneck rather than passing the Judge the full metric set.Source: §2.3 · pp. 7–9
  • Offline selection profiles 100 self-refine kernels per representative task, retains taskwise Top-20 runtime-correlated metrics, then selects 24 stable cross-task metrics above the global P75 threshold.The selected metrics are used as references when the Judge profiles generated kernels.Source: §2.3 · pp. 7–8

3 Experiments

CudaForge is evaluated on KernelBench against model, agentic, and RL-based baselines, with ablations and generalization tests examining its performance, costs, and design choices.

3.1 Benchmark and Evaluation

  • KernelBench contains 250 tasks across three difficulty levels, and evaluation reports correctness, performance, Fast1, median speedup, and 75th-percentile speedup.For iterative or multi-candidate methods, the reported kernel is the best-performing correct candidate per task; the 25-task D∗ subset preserves the full benchmark’s level proportions.Source: §3.1 · p. 9

3.2 Settings & Baselines

  • CudaForge is tested with OpenAI-o3 as both agents and ten maximum iteration rounds, alongside one-shot, self-refinement, feedback-ablation, Kevin-32B, and Agentic Baseline comparisons.The default experiments use a Quadro RTX 6000; Kevin-32B baseline results are taken from its paper.Source: §3.2 · pp. 9–10

3.3 Main Results in RTX 6000

  • 97.6% correctness, 1.677× average performance, and 70.8% Fast1 on KernelBench; on D∗, CudaForge reaches 100% correctness and 1.767× average performance.It also achieves 98% correctness and 1.776× average speedup on Levels 1–2 versus the Agentic Baseline’s 95.0% and 1.490×; on Level 3, it records 96% correctness and 1.283× average speedup. Table 1 covers the 250-task benchmark, with Agentic Baseline results limited to Levels 1–2.Source: Table 1 · p. 10, §3.3 · p. 11

3.4 Comparison with Kevin-32B on H200

  • On H200, CudaForge reaches 98.0% correctness and 1.662× performance on Levels 1–2, versus Kevin-32B’s 82.0% and 1.10×; Level 3 results are 96.0% and 1.261×.The comparison is not fully reproducible because Kevin-32B’s prompts, evaluation setup, and benchmark specifications are unavailable; CudaForge is re-evaluated on H200 to align the settings as much as possible.Source: §3.4 · pp. 12–13

3.5 API and Computation Time Cost of CudaForge

  • 26.5 minutes on one RTX 6000 and $0.3 API cost per kernel compare with the Agentic Baseline’s approximately six H100 GPU hours and $5 per kernel.CudaForge already surpasses that baseline at no more than $0.15 and ten minutes per task; the authors attribute cost efficiency to targeted hardware feedback, selective NCU metrics, and lightweight memory.Source: §3.5 · pp. 13–14

3.6 Ablation Studies

  • Using all NCU metrics lowers correctness and performance and raises cost to about 40 minutes and $1 per kernel; self-refinement and single-feedback ablations also underperform the full workflow.Self-refinement yields 1.107× speedup and 55.2% Fast1, while the correction-only Judge yields 1.222× and 58.8% Fast1; optimization-only feedback often leaves kernels unable to compile or run. The authors identify the full design’s 24-metric subset, separate Coder and Judge roles, and combined correction and optimization feedback as effective design choices.Source: §3.6 · pp. 14–15

3.7 Generalization Capability of CudaForge

  • On KernelBench’s stratified subset D*, increasing optimization rounds to 30 raises average speedup to 2.271×; gains continue from 10 to 30 rounds, but at a slower growth rate.Source: §3.7 · p. 15, Figure 7 · p. 15

4 Case study

A CrossEntropyLoss case study shows the Judge using NCU feedback to diagnose performance bottlenecks and correctness errors, then guide kernel refinement.

4 Case study

  • 1.66× to 2.42× speedup followed replacing repeated block synchronizations with warp-level shuffle reductions after 23.7% of active warps stalled on barriers.The suggested reduction lowered __syncthreads() per block from 16 to 2.Source: §4 · p. 16
  • A failed numerical check was traced to thread 0 using an uninitialized target_logit, and broadcasting the value via _shfl_sync fixed the mismatch.The numerical issue disappeared after the fix.Source: §4 · pp. 16–17
  • 3.436× to 3.762× speedup followed rounds 6 and 7, which targeted register pressure and redundant global-memory reads.The Judge recommended reducing per-thread registers and reusing logits buffered during the max pass.Source: §4 · p. 17
  • The case illustrates a workflow that first diagnoses hardware bottlenecks and then derives a corresponding optimization strategy.The figure presents representative optimization and repair rounds from a 10-round refinement process.Source: §4 · pp. 16–17

5 Supplement Observations

The authors identify “fake kernels” in CUDA-L1 results and manually check their own experimental kernels.

5 Supplement Observations

  • Some CUDA-L1 outputs reported as performant contain no custom CUDA code and instead use try-except fallbacks to PyTorch implementations.To address this evaluation concern, the authors manually checked every kernel in their experiments.Source: §5 · p. 17

6 Conclusion

CudaForge is a training-free framework that uses iterative hardware feedback to refine CUDA kernels; the conclusion reports strong KernelBench results, robustness across GPUs and base models, and low API and time costs.

6 Conclusion

  • CudaForge imitates expert iterative development by using hardware feedback for targeted refinement, and the conclusion reports the highest correctness and significant performance gains on KernelBench, robustness across GPUs and base LLMs, and scaling with refinement rounds.The authors describe its low API and time cost as making it a practical solution for automated CUDA development.Source: §6 · p. 18

A.1 Prompt for Coder in first round of generation

  • The first-round Coder prompt asks for custom CUDA replacements for PyTorch operators, permitting operator fusion and algorithmic changes such as online softmax.The prompt requests functional code rather than pseudocode.Source: Prompt for Coder in first round of… · p. 21

A.2 Prompt for Judge

  • The Judge prompt design places the role specification and output schema in the system prompt and supplies separate prompts for correction and optimization..Source: §A.2 · p. 21

Prompt for CUDA Kernel Correction

  • The Judge’s correction prompt prioritizes one critical behavioral mismatch, favoring semantic issues over micro-optimizations and style.It asks for a brief issue explanation and minimal fix hint.Source: Prompt for CUDA Kernel Correction · pp. 21–22

Prompt for CUDA Kernel Optimization

  • The Judge’s optimization prompt uses GPU specifications, reference and candidate code, and NCU metrics to identify one major bottleneck and propose one metrics-driven optimization.It prioritizes changes addressing measured conditions such as occupancy limits, register pressure, or memory coalescing.Source: Prompt for CUDA Kernel Optimization · pp. 22–23

A.3 Prompt for Coder from Round 2 to N

  • From round 2 onward, the Coder corrects or optimizes the current kernel candidate based on Judge feedback..Source: §A.3 · p. 23

Prompt for Kernel Correction

  • The correction prompt fixes compilation or runtime errors and prescribes a structured Python/CUDA extension layout while excluding test code and extra prose..Source: Prompt for Kernel Correction · p. 23

Prompt for Kernel Optimization

  • The optimization prompt applies the Judge’s suggestion while preserving correctness tolerances and the public Python API.It requests a structured CUDA extension implementation and targets latency and throughput on the specified GPU.Source: Prompt for Kernel Optimization · p. 24

B.1 Why choose NCU subset metrics?

A representative Level 2 task contrasts Judge feedback from full NCU metrics with feedback from a 24-key subset; the subset configuration guides more productive kernel refinement.

B.1 Why choose NCU subset metrics?

  • Within 10 rounds, the full-metrics Judge left kernel performance almost unchanged, whereas the 24-key subset configuration often improved performance and eventually produced a substantial speedup.The comparison is shown for Level 2 task 51.Source: Figure 9 · p. 25, §B.1 · p. 25

Response by using the full NCU metrics set

Response by using the 24 key NCU subset

  • The 24-key Judge identified the second pass over x_orig as the DRAM-bound bottleneck and proposed caching x_orig in registers to eliminate a global-memory read.The rewrite reduced traffic by about 4 MB per batch and yielded more than 30% speedup in the reported setting.Source: Response by using the 24 key NCU subset · p. 26

B.2 Top-20 NCU metrics Example

This section reports example NCU metrics ranked by their correlation with kernel runtime. The examples cover Conv2D and SpMM tasks.

B.2 Top-20 NCU metrics Example

  • The Top-20 metrics are ranked by absolute Pearson correlation with kernel execution time.Positive correlations indicate that larger metric values typically accompany longer runtimes; negative correlations indicate shorter runtimes.Source: §B.2 · p. 26
  • The Conv2D and SpMM tables show Top-20 NCU metrics and their Pearson correlations with runtime.The section ranks metrics by the absolute value of their Pearson correlation coefficient.Source: Table 6 · p. 26, §B.2 · p. 26
  • The listed examples include dram__cycles_elapsed.avg.per_second, launch__waves_per_multiprocessor, and launch__thread_count.The supplied metric entries report correlations of 0.665 784, 0.627 478, and 0.627 478, respectively.Source: §B.2 · p. 27

B.3 Key Subset of 24 NCU Metrics

This section presents a 24-metric NCU key subset produced by Algorithms 1 and 2.

B.3 Key Subset of 24 NCU Metrics

  • The 24-metric NCU key subset is the result of Algorithms 1 and 2.Table 8 is captioned as the 24-metric key subset.Source: §B.3 · p. 27, Table 8 · p. 27
  • The listed subset entries include smsp__warp_issue_stalled_branch_resolving_per_warp_active.pct and smsp__sass_average_branch_targets_threads_uniform.pct.They appear as entries 23 and 24 in the supplied list.Source: §B.3 · p. 28

C CUDA-L1

The CUDA-L1 review finds that reported high speedups often come from PyTorch-level optimizations rather than custom CUDA kernels. It also identifies fallback behavior when custom kernels fail to compile.

C CUDA-L1

  • The replication review notes that CUDA-L1 released only final generated kernels, which it examined to identify these patterns.The review describes its findings as arising from studying the released cases.Source: C CUDA-L1 · p. 28
  • Nine of CUDA-L1’s ten top speedup solutions use official PyTorch implementations rather than custom CUDA kernels.The review identifies emphasizing PyTorch-level optimizations as the most frequent issue in the provided cases.Source: C CUDA-L1 · p. 28
  • 120.3× and 64.4× reported speedups correspond to top-ranked entries that contain no CUDA kernel.The top-ranked entry is Level 2 Task 83; the second-ranked is Level-1 Task 12, Matmul with diagonal matrices.Source: C CUDA-L1 · pp. 28–29
  • Speedups clustered around 1.00, typically within ±5%, can reflect fallback to the original PyTorch operator after compilation failure.For Level-1 Task 3, fallback to torch.bmm after a custom-kernel failure yields a reported 1.006× speedup.Source: C CUDA-L1 · p. 31
  • The Level-1 Task 3 example calls torch.bmm(A, B) after a compilation failure, although replacing that operator is the task objective.The review reports that this fallback explains the 1.006× speedup.Source: C CUDA-L1 · p. 31

D Details of Benchmark

The appendix describes KernelBench’s task levels and explains the stratified subset used for detailed analysis and lower-cost experiments.

D.1 KernelBench

  • KernelBench has 270 tasks across four difficulty levels, with Levels 1–3 comprising 250 commonly evaluated tasks.The benchmark pairs each task’s PyTorch reference with fixed input-output specifications for correctness and performance validation.Source: §D.1 · p. 32
  • Levels 1, 2, and 3 cover basic operators, composite operations, and end-to-end model components, respectively.The levels test correctness, operator composition and intermediate-memory management, and efficiency under realistic workloads.Source: §D.1 · p. 32

D.2 Our stratified random subset D∗

  • The stratified subset D∗ contains 25 tasks: 10 each from Levels 1 and 2, and 5 from Level 3.It uses a fixed 10% sampling ratio per level, preserving difficulty distribution and sampling task types uniformly within each level.Source: §D.2 · pp. 32–33

Related papers

Loading 2511.01884v2…