Automatic Kernel Generation with Large Language Models: Taxonomies, Evaluations, and Cross-Platform Challenges

1. Introduction

1.1. Motivation

Large language models (LLMs) have recently become ubiquitous, being used widely for diverse daily tasks by both technical experts and general users. This rapid adoption has placed additional scrutiny on the machinery that powers these tools, especially the hardware accelerators that are most commonly implemented as neural processing units (NPUs), graphics processing units (GPUs), or tensor processing units (TPUs). NPUs, GPUs, and TPUs are specialized hardware components designed to support highly parallel AI and mathematical workloads more effectively than standard central processing units (CPUs).

To leverage the benefits of hardware accelerators, small specialized programs, called kernels, must be designed and implemented for the target hardware platform. The speed at which LLMs can run is directly related to both the hardware itself (e.g., the number of compute cores, clock speed, and memory capacity/bandwidth) and the efficiency of the kernel functions used to compute the operations. Kernel development and optimization present several challenges. First, kernels are implemented at a lower level than common high-level programming languages such as Python or C++. Second, producing highly optimized kernels requires engineers to have deep knowledge of both software development and the intricacies of the hardware platform. Third, kernel optimizations are typically required with each generation of new hardware, since the strengths and tradeoffs of each hardware cycle can change the effectiveness of software patterns. Thus, the implementation and optimization of kernels is expensive in terms of both money and time. However, the benefits of optimizing a frequently used kernel are clear: reducing the runtime of a particular kernel by even a few milliseconds can yield substantial overall speedups for LLM training and inference, leading to substantial potential cost savings.

Drawing inspiration from LLM-based coding, interest has grown in LLM-based kernel generation. If LLMs can assist humans in kernel development, or perhaps even produce kernels that surpass the performance of human-designed functions, they can provide corporations with a competitive advantage.

1.2. Relation to Existing Surveys

Despite being a young research area, LLM-based kernel generation has gained significant attention in the past one to two years. The earliest works in this field appeared in early 2025; accordingly, this survey focuses on research reported in 2025 and 2026. Although the publications are relatively recent, two surveys have already provided coarse coverage of prior work. Yu et al. [1] published an early preprint of their survey in January 2026, providing broad coverage of the field, including a general description of methods, benchmarks, and datasets, as well as a list of open challenges. A second survey by Li et al. [2], posted as a preprint in mid-2026, motivates kernel generation from a hardware perspective and provides more detailed descriptions of several papers, but notably does not cover work from 2026 in a field that is advancing rapidly. In both surveys, methods are grouped depending on whether they are agentic or training-based, which oversimplifies the research area. Moreover, both surveys focus heavily on GPU-based accelerators and provide only cursory mention of NPUs. Although both surveys discuss benchmarking and inherent issues, neither provides a detailed analysis of the benchmarking challenges that are prevalent in the field, especially from an industrial perspective.

Table 1 summarizes the topics covered and insights offered by our survey relative to the two earlier surveys. As shown in the table, the scope of our survey extends beyond summarizing the state of the art, datasets, and benchmarks, providing a more thorough and analytical view of the field.

1.3. Literature Identification and Selection Protocol

Automatic kernel generation with LLMs is a rapidly emerging area where the majority of the most recent works are arXiv preprints, technical reports, and other public resources that have not yet been published in peer-reviewed venues. Given this rapidly evolving pool of literature, we followed an SLR-inspired literature identification and screening protocol to identify works for inclusion in a more transparent and reproducible fashion. Candidate works were identified primarily through Google Scholar and arXiv searches, supplemented by backward and forward citation tracking (i.e., screening references in included papers and later papers citing them), using the two surveys mentioned in Section 1.2 as the starting point [1,2]. Some included works are also published or indexed through venues such as ACM, IEEE, Springer, or OpenReview, but these platforms were not used as independent search databases in our protocol.

Search queries combined kernel-generation terms, such as “kernel generation”, “kernel synthesis”, “kernel optimization”, “CUDA kernel”, “Triton kernel”, “AscendC”, and “NPU kernel”, with automation terms such as “large language model”, “LLM”, “agent”, “code generation”, “reinforcement learning”, and “evolutionary search”. We included papers that proposed or evaluated automatic kernel generation or optimization methods, and we retained benchmark and dataset papers when they informed the evaluation, benchmarking, or background sections. The primary search focused on works from mid-2025 through September 2026, when LLM-based automatic accelerator-kernel generation began to emerge as a distinct research topic. Earlier works were included selectively when needed for background on GPU/NPU architecture, compiler/runtime systems, benchmarking methodology, and classical kernel optimization.

For filtering and exclusions, we omitted works that focus solely on general-purpose code generation, manual kernel optimization, conventional compiler optimization without automatic generation, or non-kernel software tasks. Papers were also omitted where the primary emphasis was automatic code generation and kernel generation was presented as a supplementary, related application. Because the field is young, we adopted a permissive inclusion strategy for method papers while separating method papers from benchmark, dataset, survey, and background works.

LLM tools were used to assist candidate discovery, citation triage, and preliminary extraction, but final inclusion decisions, taxonomy assignments, and quantitative values were checked by the authors against the source papers. Table 2 summarizes the details of the SLR-inspired approach for candidate work inclusion in this survey.

1.4. Scope and Contributions

Within the context of previous surveys and our literature selection protocol, this survey provides additional organization and practical insight into the research area across hardware platforms, extending the introductions provided by prior surveys. In particular, this survey makes the following contributions beyond reviewing recent state-of-the-art work:

1.

Novel taxonomy based on the generation-time workflow. As research on LLMs and kernel generation advances rapidly, broad categorization of methods as training-based or agentic is too coarse and, more problematically, conflates orthogonal dimensions (e.g., an agentic kernel-generation method can be either training-based or training-free). In this survey, we highlight the recent trend of generation-time harnesses and define a multi-level taxonomy that captures LLM-based kernel generation methods. The first level of this taxonomy splits methods according to the type of generation-time control loop that is used, while the second level applies orthogonal tags that further categorize the method (e.g., Level 1: Search-based generation; Level 2: RL-training-based).

2.

Cross-platform emphasis. We provide a detailed case-study analysis of kernel generation on Ascend NPUs. This case-study considers the unique aspects of kernel generation on Ascend NPUs compared with GPUs, including hardware differences, datasets, and result trends that have appeared in previous literature across hardware platforms.

3.

Evaluation-integrity analysis. The field of automatic kernel generation has completed the first step of providing standardized benchmark data, but most reported results are not directly comparable because of how the benchmarks are used. The diversity of approaches, as discussed in Contribution #1, may contribute to this disparity in benchmark usage. We provide a detailed analysis of the factors that can affect benchmark results, as well as a listing of parameters that we believe subsequent work should minimally report when presenting results.

4.

Industrial focus and insights. Even if benchmark results are reported under equitable conditions, such performance comparisons may be adequate for research publications but remain insufficient for leveraging automatic kernel generation in industrial practice. We examine additional challenges and safeguards required to take automatic kernel generation systems from research ideas to industrial tools.

2. Background: Kernel Generation Tasks and Accelerator Targets

Automatic coding agents have seen widespread growth and adoption in 2026. Reports suggest that CodeX experienced exponential increases in active users in 2026, starting at under 1 million in February and surpassing 15 million by August [3]. Clearly, software developers have found that coding agents are generally useful for implementation tasks. Automatic kernel generation is a specialized form of code generation, suggesting that improvements in the latter may transfer to the former. However, several distinguishing features separate general code implementation from kernel development.

• Hardware access and abstraction. During typical software development in high-level languages such as Python, C++, and Java, engineers are not directly concerned with the low-level hardware on which the code runs. Code executes on top of the operating system, and programming languages are designed to abstract away low-level hardware details. In kernel development, this approach is not possible. Kernel code directly controls hardware registers, execution units, and cores. To extract the most from the hardware, kernels must be designed to leverage the unique strengths of the hardware.
• Execution model. For general software, program execution is typically either fully sequential or has limited parallelism implemented in a few threads on a small number of CPU cores. The execution paradigm for kernels is very different and massively parallel.
• Memory management. High-level programming languages typically provide memory management services such as automatic garbage collection [4] or heap allocation [5]. Kernel development demands manual memory management, requiring the optimization of shared memory and cache lines for best performance.
• Optimization goals. In general software development, the primary goals are to generate working code that is clean, readable, easy to maintain, and low latency. Maintenance is a significant consideration in software development, given that some critical software can remain in production for decades, e.g., [6,7]. In kernel development, raw throughput and optimal hardware utilization are the primary concerns. Readability and maintainability are much lower priorities, since kernels will typically either be migrated automatically or completely rewritten when a major accelerator architecture revision is released (e.g., NVIDIA’s roadmap follows an approximate two-year cadence for major GPU architecture releases [8]).

Despite these differences, there is still optimism and growing evidence that LLMs can make positive contributions to kernel development in a manner similar to how CodeX [9] and Claude Code [10] have affected the software development industry.

2.1. Problem Definition: Automatic Kernel Generation

The problem definition for automatic kernel generation is as follows: given a description of a desired kernel, D, an LLM model, M, an environment with which the model can interact, $\mathcal{E}$, and an evaluation harness, $\mathcal{H}$, the generation time workflow, $\mathcal{G}$ produces kernel code, K that implements the desired functionality as output. Formally,

$$K=\mathcal{G}(D,M,\mathcal{E},\mathcal{H}).$$

(1)

In this context, the environment includes tools (e.g., profiler, debugger, file editing, etc.) and agents (e.g., code modification agents, orchestrators, etc.) that may be used during the kernel development process. The evaluation harness is used to check the correctness and speed of generated kernels and may also invoke specific tools. The functional description, D can take many forms, including high-level natural language [11,12,13], unoptimized Triton, AscendC [14,15], or a baseline PyTorch implementation [16,17,18,19,20,21,22,23].

Figure 1. Conceptual description of the automatic kernel generation problem. The task that defines the kernel functionality is provided as the input to the generation time workflow. The “brain” of the process is the LLM model, but it leverages tools from the environment to profile and perform speed measurements upon the kernel. Once a final kernel is generated, it can be utilized on a target hardware platform (e.g., Ascend NPU).

Figure 1. Conceptual description of the automatic kernel generation problem. The task that defines the kernel functionality is provided as the input to the generation time workflow. The “brain” of the process is the LLM model, but it leverages tools from the environment to profile and perform speed measurements upon the kernel. Once a final kernel is generated, it can be utilized on a target hardware platform (e.g., Ascend NPU).

Currently, the most common input format is a PyTorch implementation, as it provides a concrete, runnable definition of the desired kernel functionality and can be easily profiled to determine the speed improvements provided by the LLM. On the output side, K can take a number of different forms depending on the particular task and hardware platform. Mid-level domain-specific languages (DSLs) such as Triton [24], Helion [25], TileLang [26], CuTe [27], and ThunderKittens [28] are commonly used as output formats because of their improved readability (e.g., Triton has Python-like syntax), higher-level syntax (e.g., focusing on block/tile operations rather than hardware threads), and improved cross-platform support. For additional control, the output can alternatively take the form of parallel computing platforms such as AscendC or raw CUDA C/C++.

Generally, the lower the level of the code, the more challenging the task tends to be for the LLM. For example, if the desired output of the LLM is a mid-level Triton-Ascend program, the LLM must produce only a single Pythonic file, whereas if the output format is AscendC, the LLM needs to produce five separate files: host-side (CPU) code, device-side (NPU) code, a header file, a prototype file, and a build configuration. Figure 2 provides a visual illustration of the differences in output formats.

In addition to the code abstraction level, the hardware platform also affects performance. Freely available sample code for CUDA and GPU-based platforms is more readily available and is undoubtedly harvested when performing pretraining of closed and open-source models. The relative scarcity of sample code and documentation for platforms with lower market share (e.g., Huawei Ascend) naturally leads pretrained LLM models to have less foundational knowledge of the platform, which in turn leads to lower benchmark performance.

This survey does not strictly differentiate between kernel generation and kernel optimization. Specifically, in both scenarios, an input definition, D, of the kernel is provided and the goal is to produce an output kernel implementation, K. Generally, the output implementation is expected to be correct and ideally faster than the reference implementation. Thus, this survey considers kernel optimization simply as a special case of kernel generation in which the input and output representations are the same.

2.2. Evaluation Criteria

When evaluating the quality of source code in general software development, metrics typically prioritize correctness and readability. For kernel development, the situation is different because correctness and speed are the primary concerns, while readability is not normally a primary evaluation criterion. Thus, metrics in standard benchmarks typically focus on correctness and speedups [29,30]. In addition to the metrics themselves, kernels are mathematical functions or combinations of functions. A natural question arises: How many input values must be tested to ensure that the kernel implementation is “correct”? How should boundary cases be tested? Different data types may invoke different pathways on accelerator hardware, so which data types should be included in correctness evaluations? Moreover, what levels of numerical precision are required to consider an answer “correct”? Similar to implementing mathematical functions that will run on a CPU, all of these factors must be considered for the specialized task of kernel evaluation.

Since kernel functions are defined at such a low level of abstraction, they have a number of additional aspects that must be considered during evaluation. Deep learning systems require a large number of packages. For example, on Ascend NPUs, relevant packages include torch_npu, vllm_ascend, CANN, and Triton-Ascend, and the versions of these packages can greatly affect kernel timing. Moreover, when evaluating kernels, it is typically desirable to measure the steady-state runtime of the kernel, which is achieved after performing kernel “warmup” (i.e., running the kernel for multiple passes with dummy workload inputs before timing the kernel). Warming up the kernel is desirable because the first kernel run requires hardware-specific setup, such as JIT compilation and tile calculation, which can distort the practical runtime observed during typical operation. In subsequent sections, this survey provides an in-depth analysis of the abovementioned and other practical challenges of kernel generation evaluation and proposes solutions to assist the field in moving forward.

2.3. Paper Structure

The remaining sections of the paper are structured as follows, as shown in Figure 3. In Section 3, we define our control-loop taxonomy of automatic kernel generation methods and classify related work within this framework. In Section 4, a detailed analysis of benchmarking and evaluation is provided, including current measures that are being used to produce trustworthy evaluations as well as steps that are still needed to ensure that the field is advancing appropriately. Section 5 considers cross-platform issues, focusing on a case study of kernel optimization on Ascend NPUs and analyzing trends across hardware platforms. Section 6 offers an assessment of the open problems in the field as well as a research opportunities that can be pursued to fill these gaps. Finally, Section 7 concludes the survey.

3. Automatic Kernel Generation Taxonomy based on Generation-Time Workflows

In addition to summarizing the most recent state-of-the-art work in the area of automatic kernel generation, one of the main contributions of this survey is the introduction of a new taxonomy, which uses the generation-time (or inference-time) control loop as the primary criterion for category. Other design aspects, such as learning method, representation, feedback signal, and platform are included in the taxonomy as orthogonal axes, where single or multiple relevant “tags” can be applied to each method. Previous surveys have relied on taxonomies that simply distinguish between training-free, training-based, and agent-based approaches.

3.1. Why Existing Taxonomies are Insufficient

The previous survey by Yu et al. [1] first coarsely divides prior art based on whether an agentic approach is used, defining two categories called “LLM4Kernel” and “Agent4Kernel” . The LLM4Kernel category is meant to capture works that leverage an LLM in a traditional, non-agentic fashion for kernel generation, including works that train the LLM with SFT or RL. The Agent4Kernel category is intended to capture approaches that follow a more agentic workflow, in which early works often relied on a pre-trained LLM and did not perform additional SFT/RL fine-tuning. There are two issues with this taxonomy breakdown. First, many approaches span both LLM4Kernel and Agent4LLM categories, especially as RL training is increasingly used for agentic LLM training. For instance, Yu et al. [1] classify CUDA Agent [31] in the LLM4Kernel group because it leverages RL training, but this work could also clearly be included in the Agent4LLM category. Second, the taxonomy confounds dimensions, making it difficult, or even impossible, to properly separate many modern and upcoming works. The survey by Li et al. [2] adopts a slightly different approach, grouping methods depending on whether they are single-agent, multi-agent, or training-based. Again, this categorization separates agentic and training-based approaches, even though they are not mutually exclusive.

3.2. Our Taxonomy

In this survey, we define a new taxonomy that separates methods based on the primary criterion of their generation-time workflow. This taxonomy is reflective of the current state of research in LLMs and VLMs, where agentic methods are becoming dominant. The four, largely non-overlapping categories of generation-time workflow that we define for kernel generation are: 1) direct generation, 2) structured feedback refinement, 3) search-based generation, and 4) adaptive tool-orchestrating agentic systems. These four categories are shown pictorially in Figure 4. The relative complexity of the categories generally increases from category one to four, with “direct generation” approaches essentially generating kernels by taking the output of the LLM directly, whereas the adaptive tool-orchestrating systems are typically multi-agent in nature and allow for dynamic selection of tools, memories, and expert modules.

Within each of these four primary generation-time workflow categories, there are significant variations. For example, some works will focus purely on the search aspects of the loop design and leverage a pre-trained model for performing the kernel code modifications. Other works may consider both the interaction of RL training for kernel generation model in the context of a search loop based system design. To accommodate such diversity, we define a set of orthogonal axes that can be used to describe each method more precisely. These four orthogonal design axes are shown in Figure 4. For example, in the “learning” axis, we include the sub-categories of “None”, “SFT”, “RL”, “DPO”, etc. These sub-categories are meant to be used as tags where one or more can be used, as appropriate. An advantageous aspect of our taxonomy design is that it can be extended easily in the future as new techniques are developed by adding additional inference-time control loop categories or new tags along the orthogonal axes.

In the following subsections, we first define the rules for assigning methods to the taxonomy categories and subsequently define each category more precisely with corresponding example works from the literature.

3.3. Taxonomy Category Assignment

In order to ensure that the taxonomy is objective and reproducible, we provide a set of rules in the form of a binary decision tree. This series of three simple yes/no questions can be used to map existing methods into the appropriate taxonomy category. The first question identifies whether a search state is maintained, which is a clear indicator of membership to Category #3. For methods that do not perform search-based generation, the remaining questions identify whether a dynamic, agentic workflow, a fixed revision loop, or a single forward pass are employed.

Table 3 provides further guidance regarding the rules/questions used to assign methods to the taxonomy categories and is complimentary to Figure 5. The table provides a particularly useful column to help distinguish between borderline cases. A key guiding principle when assigning methods to the taxonomy is that when mulitple properties are present, the approach is classified using the mechanism that determines the next generation action. For example, if a method uses profiler feedback but the next candidate is chosen by a population/tree search procedure, it is assigned to the Search-Based Generation category. If the method uses named agents, but the algorithmic steps are always generator → evaluator → refiner, it belongs to the Structure Feedback category, not adaptive orchestration. If it uses SFT/RL-trained models but inference is a single pass, it is still considered a member of the Direct Generation category because training is an orthogonal learning axis.

3.4. Category #1: Direct Generation

The first category, “direct generation”, is the most basic approach to kernel generation. In this case, the generated kernel is commonly the direct output of a single model computed in one forward or pipeline pass. In the simplest setting, given an input description of the desired kernel functionality, D (e.g., textual description, reference PyTorch implementation, etc.), the model produces the desired kernel output, K, in a single turn. This category may also be referred to as “bare model” or “harness-free”. Historically, papers that focus on direct generation have highlighted unique training strategies for the task of kernel generation and optimization. More recently, the direct generation approach is used more as a baseline to measure the improvement of leveraging a more advanced generation-time workflow while the model is held fixed.

Figure 6. Conceptual diagram of Category #1: Direct Generation. In this category, kernels are generated in a single forward pass. The forward pass may consist of multiple steps, but no attempts are made to refine the solution based on previous attempts.

Figure 6. Conceptual diagram of Category #1: Direct Generation. In this category, kernels are generated in a single forward pass. The forward pass may consist of multiple steps, but no attempts are made to refine the solution based on previous attempts.

Most works that fall in the direct generation category focus on the data collection and training strategies. ConCur [16] presents a data generation pipeline that produces high-quality reasoning traces that are then used in LoRA fine-tuning a Qwen model which shows improved single-step kernel generation. AutoTriton [17] focuses on a sophisticated PyTorch to Triton LLM distillation data generation pipeline as well as its SFT and GRPO-based training approach, but ultimately performs single-shot Triton kernel generation at inference. TritonRL adopts a similar approach, but also introduces a new hierarchical reward which improves credit assignment during long-sequence generation [32]. AscendKernelGen devises an extensive data collection pipeline that leverages an ensemble of LLMs to create a new dataset targeting Ascend NPUs, called NPUKernelBench [33]. The paper also shows an SFT + RL training pipeline that utilizes the new dataset and demonstrates that the training procedure yields an improved model that outperforms the comparable baseline for direct NPU kernel generation tasks [33].

Other work develops an advanced training pipeline consisting of SFT with data augmentation, self-supervised learning, and contrastive RL and investigates the contributions of each training phase [18]. CUDA-L2 [34] builds upon the previous work and employs a specialized RL stage that focuses on half-precision general matrix multiplication (GEMM) to show significant gains over the highly optimized cuBLAS library.

Although most of the direct generation methods can be considered as baseline, “harness-free” methods, that is not a strict requirement. AscendCraft [35], for instance, considers the more challenging problem of AscendC kernel generation, but adopts a single forward pass pipeline that first generates an intermediate-level DSL and subsequently performs a constrained transcompilation process that operates upon code slices. Although some level of feedback to the LLMs is provided in order to correct for compiler errors during lowering, AscendCraft is an example of a multi-stage direct generation pipeline for AscendC kernel generation. The work in QiMeng-Attention [36] similarly adopts a multi-stage generation pipeline that uses an intermediate “thinking language” representation to allow an LLM to optimize complex attention operators on GPU.

Table 4. Comparison of Category #1 direct generation kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
ConCur [16]	SFT (LoRA)	PyTorch	CUDA	N
AutoTriton [17]	SFT, RL	PyTorch	Triton	N
TritonRL [32]	SFT, RL	PyTorch	Triton	N
CUDA-L1 [18]	SFT, RL (contrastive)	PyTorch	CUDA	N
CUDA-L2 [34]	Pretraining, RL (contrastive)	HGEMM task specification	CUDA	N
Ascend-KernelGen [33]	SFT, DPO	Task specification	AscendC	A
QiMeng-Attention [36]	None	Attention algorithm specification	CUDA/CuTe	N
Ascend Craft [35]	None	PyTorch	AscendC	A

Table 4. Comparison of Category #1 direct generation kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
ConCur [16]	SFT (LoRA)	PyTorch	CUDA	N
AutoTriton [17]	SFT, RL	PyTorch	Triton	N
TritonRL [32]	SFT, RL	PyTorch	Triton	N
CUDA-L1 [18]	SFT, RL (contrastive)	PyTorch	CUDA	N
CUDA-L2 [34]	Pretraining, RL (contrastive)	HGEMM task specification	CUDA	N
Ascend-KernelGen [33]	SFT, DPO	Task specification	AscendC	A
QiMeng-Attention [36]	None	Attention algorithm specification	CUDA/CuTe	N
Ascend Craft [35]	None	PyTorch	AscendC	A

3.5. Category #2: Refinement via Structured Feedback

The second category in our taxonomy takes inspiration from how human engineers perform the task of kernel generation and optimization. In particular, humans pursue an iterative approach, whereby they implement an initial solution, measure its performance, identify the bottlenecks, and subsequently refine the solution. Human kernel engineers may go through tens of rounds of refinement before arriving at an optimized kernel. Taking inspiration from this mindset, researchers have adapted this approach to kernel generation to allow a model to iteratively refine an initial solution across multiple rounds. In this multi-turn setup, feedback, which is provided in the form of compilation/profiler outputs, timing analyses, etc., is typically incorporated using deterministic rules, which gives the model hints about computational bottlenecks that require attention. It is important to stress that in this class of solutions, the feedback is provided in a deterministic manner, rather than a fully agentic set-up.

Figure 7. Conceptual diagram of Category #2: Refinement via Structured Feedback. In this category, kernels are refined across multiple turns, or steps, using feedback provided by tools such as profilers and compilers. The exact approach through which the refinement is performed is a design choice of the approach. Some methods may refactor the entire kernel during each step; whereas, others may focus on a specific code fragment during each iteration. This category allows flexibility for either such strategy.

Figure 7. Conceptual diagram of Category #2: Refinement via Structured Feedback. In this category, kernels are refined across multiple turns, or steps, using feedback provided by tools such as profilers and compilers. The exact approach through which the refinement is performed is a design choice of the approach. Some methods may refactor the entire kernel during each step; whereas, others may focus on a specific code fragment during each iteration. This category allows flexibility for either such strategy.

One of the earliest works to incorporate feedback within a multi-turn RL-trained framework was Kevin, which included execution results, timing profiles, and the previous kernel code during each step of multi-turn CUDA kernel generation [37]. CUDA Agent [31] follows a similar overall approach, but scales up the training (e.g., context length, synthesized training set, model size) and explicitly defines a skill-integrated environment. Dr. Kernel also used an generation-time loop with a single LLM that performed multi-turn generation with profiling feedback [30]. Highlights from this work include the proposal of an improved RL loss, Turn-level Reinforce-Leave-One-Out (TRLOO), which removes a bias issue inherent with GRPO, and KernelGym, a robust GPU-based sandbox environment that runs generated kernels and includes correctness checks, profiling, and reward-hacking checks. daVinci-kernel [38] extends Dr. Kernel to a feedback-style multi-agent workflow while still focusing on SFT and RL training aspects. The use of memory banks was also explored for the iterative generation of AscendC kernels, where RL was used to update the memory retrieval policy while employing a fixed LLM for kernel synthesis [39]. In [40], a novel approach introduces a Macro Thinking Micro Coding (MTMC) method whereby a lightweight RL-trained policy model proposes semantic optimization actions and a general-purpose LLM realizes these changes via code updates. This work does not include any compilation/profile feedback, but it still performs an iterative approach to kernel refinement.

Although the aforementioned works all incorporate some aspects of SFT/RL fine-tuning, that is not a necessary condition for this taxonomy category. For example, PEAK employs an iterative performance workflow that includes correctness and performance indicators [41]. It makes use of general-purpose LLMs by specifying the kernel generation operations in natural language, making them easier for a general LLM to interpret. Other work has demonstrated the use of a finite state machine (FSM) to control the overall workflow with compilation/runtime information feedback being used for Triton PyTorch ATen kernel generation on emerging ML ASIC accelerators [42]. Astra [43], GEAK [11], PRAGMA [44] and KForge [45] all adopt a feedback-based multi-agent approaches that divides kernel optimization across a set of tasks that are handled by specific agents (e.g., planning, coding, testing, and profiling agents). Even though agents are used, the optimization loops are fixed, without a central orchestrator that controls the use of tools. For this reason, these works fit more naturally in this taxonomy category than the adaptive tool-orchestrating agent class. CUDAForge [46] also takes a multi-agent approach that consists of a Coder Agent and a Judge Agent, whereby the judge uses profiling tools to provide bottleneck information to the Coder for the next round.

Other work that is included in this category has taken more of an optimization-oriented approach, where an initial kernel in a target language is provided as input and the system needs to improve performance. SwizzlePerf is an example of such work and it uses swizzling (i.e., remapping the order in which data, tiles, and memory are accessed) to achieve the optimization goal [47]. During each iteration, it creates a bottleneck report based on the compilation/validation status of the current kernel alongside the output of AMD’s rocprof profiler [48]. These outputs are used to define the swizzling strategy for the next step. TritonForge is another optimization-oriented work for Triton kernel generation [49]. It uses Gemini-2.5-Pro off-the-shelf in a straightforward, pre-defined loop structure that leverages feedback from NVIDIA’s Nsight compute profiler.

Table 5. Comparison of Category #2 refinement via structured feedback kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
Kevin [37]	RL (single- and multi-turn)	PyTorch	CUDA	N
Dr. Kernel [30]	SFT, RL (multi-turn)	PyTorch	Triton	N
QiMeng-Kernel [40]	RL (macro thinking)	PyTorch	Triton	N
PEAK [41]	None	Kernel context	CUDA, HIP, HLSL	N
TritorX [42]	None	ATen operator signature	Triton PyTorch ATen	META MTIA
SwizzlePerf [47]	None	Triton (unoptimized)	Triton	AMD GPU
TritonForge [49]	None	Triton (unoptimized)	Triton	N
daVinci-Kernel [38]	SFT, RL	PyTorch	Triton	N
CUDA Agent [31]	PPO (single turn), RFT (rejection fine-tuning), agentic PPO	PyTorch	CUDA with PyTorch	N
CUDA Forge [46]	None	PyTorch	PyTorch with CUDA extensions	N
GEAK [11]	None	Natural language, reference code	Triton	AMD GPU
KForge [45]	None	PyTorch	CUDA, Triton, CuTe, HIP, SYCL, Metal	N, Intel GPU, AMD GPU
PRAGMA [44]	None	PyTorch	Triton (GPU), C++ (CPU)	N, Intel CPU
EvoKernel [39]	RL (for memory retrieval)	PyTorch	AscendC (NPU), CUDA (GPU)	A, N
Astra [43]	None	CUDA (unoptimized)	CUDA	N

Table 5. Comparison of Category #2 refinement via structured feedback kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
Kevin [37]	RL (single- and multi-turn)	PyTorch	CUDA	N
Dr. Kernel [30]	SFT, RL (multi-turn)	PyTorch	Triton	N
QiMeng-Kernel [40]	RL (macro thinking)	PyTorch	Triton	N
PEAK [41]	None	Kernel context	CUDA, HIP, HLSL	N
TritorX [42]	None	ATen operator signature	Triton PyTorch ATen	META MTIA
SwizzlePerf [47]	None	Triton (unoptimized)	Triton	AMD GPU
TritonForge [49]	None	Triton (unoptimized)	Triton	N
daVinci-Kernel [38]	SFT, RL	PyTorch	Triton	N
CUDA Agent [31]	PPO (single turn), RFT (rejection fine-tuning), agentic PPO	PyTorch	CUDA with PyTorch	N
CUDA Forge [46]	None	PyTorch	PyTorch with CUDA extensions	N
GEAK [11]	None	Natural language, reference code	Triton	AMD GPU
KForge [45]	None	PyTorch	CUDA, Triton, CuTe, HIP, SYCL, Metal	N, Intel GPU, AMD GPU
PRAGMA [44]	None	PyTorch	Triton (GPU), C++ (CPU)	N, Intel CPU
EvoKernel [39]	RL (for memory retrieval)	PyTorch	AscendC (NPU), CUDA (GPU)	A, N
Astra [43]	None	CUDA (unoptimized)	CUDA	N

3.6. Category #3: Search-Based Generation

Another approach to kernel generation frames the task as a search process in a high-dimensional space. The goal is therefore to have a model that can effectively perform modifications/updates to the kernel code as well as an efficient search strategy to efficiently explore the most promising regions of the solution space. Since efficient search in high-dimensional solution spaces is a well-studied, long-standing problem, works using this method can leverage general-purpose solutions for the specialized task of kernel generation. For instance, a number of works maintain a population of candidate kernels and perform local perturbations upon members of the population, in a similar fashion to long-standing simulated annealing [50] and genetic-algorithm-style [51] approaches.

Figure 8. Conceptual diagram of Category #3: Search-Based Generation. The distinguishing feature of this category is that a population of kernel solutions is maintained and refined across step. As shown in the “Population/Search Space” block, the set of solutions may be organized in various ways, such as in a tree or across multiple islands of populations. These approaches typically require a mechanism to select which element from the set of current kernels will be refined, which can optionally be performed by an LLM, as shown in the “Search Controller” block. The refinement to the kernel solution itself is performed by the “LLM Generator”.

Figure 8. Conceptual diagram of Category #3: Search-Based Generation. The distinguishing feature of this category is that a population of kernel solutions is maintained and refined across step. As shown in the “Population/Search Space” block, the set of solutions may be organized in various ways, such as in a tree or across multiple islands of populations. These approaches typically require a mechanism to select which element from the set of current kernels will be refined, which can optionally be performed by an LLM, as shown in the “Search Controller” block. The refinement to the kernel solution itself is performed by the “LLM Generator”.

Perhaps the most straightforward search-based approaches that share the most similarities to structured feedback methods follow a commit-reset mindset. AutoKernel uses a simple search loop with a fixed number of iterations that performs the following steps: 1) write a new candidate, 2) benchmark the candidate, and 3) retain code improvements while discarding regressions [52]. Although this solution only maintains a single kernel implementation, it still differs from structured feedback methods due to its explicit commit-reset search strategy. The approach in CUDA-LLM shares conceptually similarities with the structured feedback category, and contains an iterative search and reinforcement loop that prompts an LLM to generate N kernels [12]. However, unlike standard test-time scaling (TTS) approaches, these kernels are evaluated against each other within the loop and the best kernel becomes the starting point in the following iteration.

Increasing in complexity, QiMeng-GEMM uses a beam search over optimization meta-prompts that consist of universal templates for general optimization techniques as well as platform-specific details where all but the top-k solutions are pruned at the end of the loop [53]. Dr. Triton also employs a more complex test-time search strategy which subdivides complex kernels into subsequences and performs a systematic search over fusion methods [54]. KernelBand formalizes kernel optimization as a multi-armed bandit (MAB) problem and uses runtime clustering as well as hardware profiling to prune weak actions and select those that are more promising. This setup provides additional structure and guidance to the LLM when it performs code modifications. In Lange et al. [55], a two-stage procedure is followed to generate CUDA kernels whereby PyTorch is first translated to an initial CUDA implementation which is optimized in the second stage. The first stage takes a structured feedback approach, but the second stage uses a population-based evolutionary search with top-k filtering. A two-stage procedure is also adopted in AscendOptimizer, which considers the more challenging case of generating AscendC code and requires both host-side tiling and device-side kernel optimizations [14]. The method repeatedly rewrites kernels in stage 1, filtering out slower options while stage 2 adopts a mutation-based search to optimize the host-side tiling.

One popular methodology in this search-oriented category is to leverage tree structures to organize the search (e.g., Monte Carlo Tree Search) within the generation-time harness and perform more targeted code optimizations. AdaExplore [56] and STARK [20], for example, maintain a collection of candidate kernels organized in a tree structure, where each node corresponds to a kernel implementation. AdaExplore leverages a failure-driven skill memory and oscillates between small and large code updates to arrive at optimized kernels; whereas, STARK uses an epsilon-greedy node selection policy alongside separate plan, code, and debug agents. QiMeng-TensorOp [13] implements an LLM-assisted MCTS auto-tuning stage in its operator generation pipeline. It uses a four-step process of node selection, expansion, simulation, and backpropagation to uncover optimization opportunities [13]. In ReGraphT, the reference is provided as sequential C/C++ code and during inference Monte Carlo Graph Search selects a sequence of optimization methods to generate the desired CUDA code [57].

In their approach, K-search combines a search tree with a world model for CUDA kernel generation [58]. The elements in the tree are either closed nodes, representing completed kernels, or open nodes, corresponding to hypotheses for code changes. K-search selects an open node from the tree, performs the corresponding local code refinement, and uses a world model to update the search tree based on the result of the code modification via insert, update, and prune operations. Other works do not use a world model per se, but employ a memory or knowledge base [23] of successes, failures, or both in order to support planning during kernel optimization. KernelEvolve demonstrates impressive kernel generation that generalizes across programming abstractions and hardware by leveraging a search graph with nodes corresponding to kernel implementations and edges denoting code transformations [59]. The generation flow is controlled by an iterative search loop that performs node selection, kernel code modifications, and scoring using a set of specialized tools or agents.

Another direction within the search-based generation class is structured around evolutionary search. These approaches are based on classical black-box optimization methods such as genetic algorithms and beam search and treat candidate kernels as populations, repeatedly generating kernel variants through mutations and crossovers. For each variant that is created, its correctness and performance are measured and only the best candidates are propagated and recombined in future iterations. EVOEngineer is an earlier work along this direction, which provided additional structure to the problem by proposing a LLM-based code evolution framework for kernel optimization [60]. The paper proposed two primary components in the EvoEngineer framework: 1) Traverse techniques, which includes the general optimization strategy and its effect on the prompt design and 2) Population management, which defines how candidate kernels are selected and evolved across generations. cuPilot is an evolutionary approach which argues that performing crossovers directly on the raw kernel code is ineffective and instead proposes a “strategy” as an intermediary representation. Crossovers are performed upon the strategies and a strategy-to-kernel translation is used to apply the crossover to the kernel code itself [61]. GPU Kernel Scientist considered an evolutionary approach for AMD GPU kernel optimization that included an evolutionary selector, experimental designer, and kernel writer in its fixed search loop [62].

In addition to methods consisting of single populations of solutions, so-called “multi-island” approaches extend the setup to multiple, mostly independent populations (called “islands”) evolving in parallel. Such a setup can be desirable when searching over complex solution spaces, as they allow the set of populations to explore a wider range of optimization strategies, where each island has its own history and specialization. Kernel Foundry is an example of such a system, which provides detailed diagnosis feedback to the islands to drive each population forward, as well as an experience library, which provides a collection of reusable optimization and repair hints that are shared across islands [63]. Whereas Kernel Foundry demonstrated a multi-island evolutionary approach that employed closed and open-source pre-trained models, Kernel-Smith derived a training recipe in which the LLM-based code modifier (responsible for the code mutations and recombinations for the candidates across populations) was improved to operate more effectively inside of an evolutionary loop [64].

Recent works have also considered multiple optimization schemes in search of the most effective method. For example, Nagaitsev et al. [19] showed that a branching-search evolutionary strategy that prioritizes exploitation outperforms an OpenEvolve-based [65] solution which emphasizes both exploration and exploitation, especially for Triton kernel generation. Interestingly, the OpenEvolve-based search yielded better results for the CUDA backend, which motivates the need for additional studies to explore the generality of these findings.

Table 6. Comparison of Category #3 search-based kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU.

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
QiMeng-TensorOp [13]	None	Natural language description	C sketch + assembly, CUDA/CuTE	RISC-V CPU, ARM CPU, NVIDIA GPU
AdaExplore [56]	None	PyTorch	Triton	N
AutoKernel [52]	None	PyTorch	Triton, CUDA	N
CUDA-LLM [12]	None	Natural language description	CUDA	N
cuPilot [61]	None	PyTorch	CUDA	N
Dr. Triton [54]	SFT, curriculum RL, RLVR	PyTorch	Triton	N
EvoEngineer [60]	None	PyTorch + C++ CUDA (unoptimized)	CUDA	N
GPU Kernel Scientist [62]	None	PyTorch	HIP	AMD GPU
K-Search [58]	ICL (world model)	PyTorch	CUDA	N
Kernel Foundry [63]	None	PyTorch	Triton	N
Kernel Smith [64]	SFT, RL	PyTorch, CUDA	Triton, MACA	N, MetaX GPU
KernelBand [15]	None	Triton (unoptimized)	Triton	N
KernelEvolve [59]	None	Kernel specification or reference implementation (i.e., PyTorch)	Triton (primarily), also CuTe	N, AMD GPU, META MTIA
PIKE [19]	None	PyTorch	Triton / CUDA	N
ReGraphT [57]	None	Sequential C/C++	CUDA	N
STARK [20]	None	PyTorch	PyTorch with embedded CUDA	N
Robust-KBench [55]	None	PyTorch	CUDA	N
AscendOptimizer [14]	None	AscendC (unoptimized)	AscendC	A
QiMeng-GEMM [53]	None	GEMM spec.	CUDA, C	N, RISC-V CPU
KernelBlaster [23]	None	PyTorch	CUDA	N

Table 6. Comparison of Category #3 search-based kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU.

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
QiMeng-TensorOp [13]	None	Natural language description	C sketch + assembly, CUDA/CuTE	RISC-V CPU, ARM CPU, NVIDIA GPU
AdaExplore [56]	None	PyTorch	Triton	N
AutoKernel [52]	None	PyTorch	Triton, CUDA	N
CUDA-LLM [12]	None	Natural language description	CUDA	N
cuPilot [61]	None	PyTorch	CUDA	N
Dr. Triton [54]	SFT, curriculum RL, RLVR	PyTorch	Triton	N
EvoEngineer [60]	None	PyTorch + C++ CUDA (unoptimized)	CUDA	N
GPU Kernel Scientist [62]	None	PyTorch	HIP	AMD GPU
K-Search [58]	ICL (world model)	PyTorch	CUDA	N
Kernel Foundry [63]	None	PyTorch	Triton	N
Kernel Smith [64]	SFT, RL	PyTorch, CUDA	Triton, MACA	N, MetaX GPU
KernelBand [15]	None	Triton (unoptimized)	Triton	N
KernelEvolve [59]	None	Kernel specification or reference implementation (i.e., PyTorch)	Triton (primarily), also CuTe	N, AMD GPU, META MTIA
PIKE [19]	None	PyTorch	Triton / CUDA	N
ReGraphT [57]	None	Sequential C/C++	CUDA	N
STARK [20]	None	PyTorch	PyTorch with embedded CUDA	N
Robust-KBench [55]	None	PyTorch	CUDA	N
AscendOptimizer [14]	None	AscendC (unoptimized)	AscendC	A
QiMeng-GEMM [53]	None	GEMM spec.	CUDA, C	N, RISC-V CPU
KernelBlaster [23]	None	PyTorch	CUDA	N

3.7. Category #4: Adaptive Tool-Orchestrating Agents

The final category of approaches considers a truly agentic workflow. In this case, the workflow is highly dynamic and adaptive—rather than following a fixed series of steps, the workflow adapts based on the current status and the result of the previous step. Moreover, it maintains a set of specialized tools and agents that can be employed to achieve the generation and optimization goals. For example, if a Coding Agent produces a kernel that repeats a previous solution, does not seem promising, or does not compile, the agentic workflow will identify that there is no need to profile this solution and instead move on to other tasks, such as calling a Designer Agent to perform an algorithmic re-design. As shown in Figure 9, one possible architecture for an Adaptive Tool-Orchestrating Agentic system is to include a central orchestrator agent that plans an optimization strategy, receives the output from all agents and tools that are invoked, and updates the strategy and next steps based on the outputs it receives. However, this architecture can be considered as just one example and is not a requirement for this category.

A prototypical example of a system that adopts the adaptive tool-orchestrating agents approach is AKG Kernel Agent [66]. In their approach, a conductor serves as the central orchestrator that dynamically leverages designer, coder, and verifier agents along with a set of tools (e.g., compile, accuracy, and profile) in its iterative loop. This type of dynamic workflow allows the conductor to invoke agents and tools based on the current state and error analysis, rather than fixed pipelines that follow a predetermined series of steps regardless of outcomes. KernelFalcon [67] adopts a similar mindset of employing a central orchestrator, but in this case the orchestrator is composed of deterministic code. KernelFalcon also maintains a persistent memory that stores artifacts such as logs, prompts, and error traces for debugging and resumption.

Figure 9. Conceptual diagram of Category #4: Adaptive Tool-Orchestrating Agents. The distinguishing feature of this category is the dynamic, adaptive nature of the workflow, whereby the tools and agents that are deployed depend upon the current state of kernel development and the feedback that is received. In this conceptual example, a central “Orchestrator Agent” controls the overall flow of the system, but other architectures are possible.

Figure 9. Conceptual diagram of Category #4: Adaptive Tool-Orchestrating Agents. The distinguishing feature of this category is the dynamic, adaptive nature of the workflow, whereby the tools and agents that are deployed depend upon the current state of kernel development and the feedback that is received. In this conceptual example, a central “Orchestrator Agent” controls the overall flow of the system, but other architectures are possible.

Table 7. Comparison of Category #4 adaptive tool-orchestrating agent kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU, C = CPU.

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
AKG Kernel Agent [66]	None	PyTorch, MindSpore, Triton, Numpy	Triton, TileLang, CUDA-C, CPP, AscendC	A, N, C
EGG [21]	None	PyTorch	Triton	N
AKO [68]	None	CUDA, Triton, TileLang, CuTe, HIP	Triton, CUDA, C++, TileLang, CuTe, PyTorch, HIP	N
KernelSkill [22]	None	PyTorch	CUDA	N
KernelFalcon [67]	None	PyTorch	Triton	N

Table 7. Comparison of Category #4 adaptive tool-orchestrating agent kernel-generation methods. A = Ascend NPU, N = NVIDIA GPU, C = CPU.

Method /Paper	Learning	Input Repr.	Output Repr.	Platform
AKG Kernel Agent [66]	None	PyTorch, MindSpore, Triton, Numpy	Triton, TileLang, CUDA-C, CPP, AscendC	A, N, C
EGG [21]	None	PyTorch	Triton	N
AKO [68]	None	CUDA, Triton, TileLang, CuTe, HIP	Triton, CUDA, C++, TileLang, CuTe, PyTorch, HIP	N
KernelSkill [22]	None	PyTorch	CUDA	N
KernelFalcon [67]	None	PyTorch	Triton	N

The recently proposed EGG [21] does not contain a central orchestrator, but it uses a multi-agent workflow that adapts based on agent outputs (e.g., after the code agent performs an update, the result is forwarded to either the debug agent or the profile agent, depending on the result of a runtime analysis). EGG also divides the overall optimization process into two main stages, independently focused on algorithmic and hardware-specific optimizations. KernelSkill is similar to EGG in the sense that it introduces a multi-agent workflow without a central orchestrator, but it consists of a complex set of LLM and deterministic agents, including LLM-based feature extractor, planner, diagnoser, optimizer, and repairer [22]. KernelSkill also leverages a dual-memory architecture consisting of a long-term memory that maintains reusable expert skills and a short-term memory that prevents repetitive backtracking. AKO [68] is another example of an adaptive agentic approach that leverages the off-the-shelf Claude Code coding agent. This work proposes a thin, task-specific harness that provides scoring, reference material, and orchestration capabilities that are typically missing in a generic coding agent for kernel optimization. In its most advanced setting, the skills and tooling in AKO’s task-specific harness can co-evolve with the generated kernels.

To conclude this section, Figure 10 provides a timeline that shows the publication dates of a sampling of the most influential work covered in this section. There are two interesting findings to note: 1) The amount of work that has been completed in such a short period of time on the topic of kernel generation speaks to the significant importance and interest in this research area. The timeline in Figure 10 covers a year and a half, which corresponds to the time window when almost all LLM-based kernel generation methods were published. 2) Research is active and distributed fairly evenly across the timeline for all four categories in our taxonomy. Categories #1 and #2 allow researchers to better isolate the contributions of LLM-training on the model’s kernel generation capabilities. Categories #3 and #4 construct more elaborate generation-time harnesses around the model, generally permitting superior performance, but also making it more challenging to precisely monitor gains resulting from improvements to the underlying model itself. Given the number of papers cited from each category, it is clear that Category #2 and #3 have been the most active directions. Our intuition is that these directions will continue to show activity, while Category #4 will grow as more complex agent systems mature.

4. Fair, Trustworthy, and Reproducible Evaluations

This survey has highlighted the relative youth of the field of automatic kernel generation and optimization. As with any new research area, kernel generation has experienced growth challenges, and the primary issues are related to organization and fair evaluation. We hope that our novel taxonomy helps to better organize methods with respect to what is most important in the real world: inference-time performance. With respect to the second issue of fair evaluations, the field has made strides in this direction in terms of releasing standardized benchmarks; however, with a wide diversity of methods and evaluation settings, it remains quite challenging to compare quantitative results directly to gain meaningful insights regarding which approach is leading the field. In this section, the current evaluation configurations are outlined, including the challenges that are faced. We then discuss ways to standardize the reporting of results, in an attempt to make result comparisons more feasible.

4.1. Evaluation Environments: Benchmarks, Harnesses, Integrity Policies, and Metrics

As shown in Section 3, automatic kernel generation has experienced significant growth in the past 1-2 years. Most of the reported papers provide quantitative experimental results and claim to be advancing state-of-the-art (SoTA) boundaries. However, one critical aspect of a fair evaluation that must be emphasized is that reported results on kernel generation are only meaningful when the benchmark, evaluation harness, evaluation integrity policy, and evaluation metrics are all clearly specified. In most academic papers, the benchmark and metrics are made relatively clear, but other aspects can only be inferred from an official code release, if at all.

To ensure common terminology, Table 8 introduces the components that are required in a full and fair pipeline for evaluating an automatic kernel generation system. There is some overlap between these four components. For instance, the KernelBench benchmark [29] includes not only the data samples themselves but also a harness for running the evaluation. However, it is completely possible to decouple these two parts and use either separately — the data samples can be used with a custom harness or the KernelBench harness could be used with different data. Even though benchmark releases often include both components, logically they are independent. There is also overlap in the sense that the Evaluation Integrity Policy is often implemented inside the of the harness, but it is a somewhat of a modularized component. Keeping these two aspects separate is useful, because two papers may use similar harness mechanics but have very different anti-hack detection mechanisms.

In the following, each of the aforementioned evaluation components will be covered in more detail.

4.2. Benchmark Data

The benchmark or dataset is the one aspect of a fair evaluation that almost every academic paper will report. The benchmark data defines the problem under investigation, describing the task and asking for a specific solution. As with most LLM-oriented datasets, a critical aspect of the dataset is how the question and LLM prompt are worded, since this explicitly defines the task. Unlike other LLM tasks, such as QA [73], code generation [74], mathematical problem solving [75], grammar error correction [76], kernel generation differs in that a “ground-truth” answer is not provided. The dataset may include in-context learning examples to aid the model, as well as a reference implementation. Notably, the reference implementation is not the expected output of the model. Rather, it is an initial baseline specification that defines the mathematical operation the LLM is attempting to implement or optimize. In some work, the reference is not a source code implementation at all, but rather a natural language description of the kernel function to be implemented [11,12,13]. Figure 11 shows a pictorial representation from the KernelBench dataset, which consists of five components: 1) The PyTorch reference implementation. 2) A function called get_inputs() to generate the test input values. 3) A function called get_init_inputs() that sets initialization values like the shapes of the test inputs. 4) Metadata that describes the task. 5) The LLM prompt, which defines the problem or task that the LLM is supposed to solve and provides in-context examples.

Another notable difference between kernel generation and optimization compared to other tasks is the close dependence on the hardware platform. For other tasks, such as computer vision, QA, etc., the hardware platform is completely unrelated to the solution. For example, it does not matter whether a new approach to QA is running on Ascend NPU or NVIDIA GPU. However, for kernel generation, the hardware platform is a critical component. If the model is expected to generate kernels that can run on Ascend NPU, then the generated kernels must be written in either AscendC or a compatible DSL (e.g., Triton-Ascend). Note that, in its original form, the KernelBench dataset [29] was intended for NVIDIA GPUs and prompted the model to generate CUDA kernels. It is possible to modify the KernelBench samples to be applicable to other hardware. For example, to adapt KernelBench for generating Triton-Ascend kernels on Ascend NPU, the prompt could be adjusted to inform the model: “You write custom Triton-Ascend kernels to replace PyTorch operators...” and “Given this PyTorch architecture, generate ModelNew with custom Triton-Ascend kernels”. Further, within the reference code and in-context examples, CUDA-specific commands require modification, e.g., $a=torch.randn(1,128).cuda\left(\right)\to a=torch.randn(1,128).to(“npu”)$. Such dataset modifications are minor: the core set of tasks or kernel functions that the LLM must generate has not changed, only the requested output format and provided code blocks are modified. However, as a community, to ensure result reproducibility, any such modifications to the dataset (e.g., reconfiguring for different hardware platforms) should be reported and the modified data should be released.

Fair Evaluation Challenge #1: Benchmark datasets are often modified by independent researchers for various reasons, such as adaptation to a different hardware platform or DSL. Such modifications may make evaluations inconsistent and incomparable. To ensure fair evaluations, authors should clearly note such modifications and release the modified benchmark.

More recently, datasets such as MultiKernelBench [69] have been released that accommodate multiple hardware platforms and output formats, which also help alleviate this challenge.

4.3. Evaluation Harness

Performing the evaluation for kernel generation is not a straightforward process. It requires significant care to perform properly. Evaluation consists of many steps, including loading the reference and LLM-generated implementations, completing warm-up runs on the kernels before timing to exclude one-time overheads (e.g., kernel loading, JIT compilation, cache effects, etc.), performing precise timing of the kernels, checking the correctness of the output to specific numerical precision, and profiling the kernels. These tasks are typically scheduled in parallel across multiple accelerators for efficiency. From this description, it is clear that there are a number of parameter settings involved in the harness, such as compilation settings, number of kernel warm-up attempts, number of timing runs per kernel (if taking an average), numerical precision settings (may be universal or data type specific), timeout duration, number of retry attempts if a failure occurs, etc. These settings will all affect the quantitative results that are reported.

Fortunately, most benchmarks are released with evaluation harness code or other APIs for performing a standardized evaluation. However, since this application is so closely tied to industrial considerations, companies will often choose to utilize their own harness during evaluation. For example, on Ascend NPUs there is an official CANNBot execution harness that is meant to assist in the creation, debugging, and testing of AscendC kernels [77,78]. The harnesses from KernelBench and CANNBot differ in terms of the correctness check numerical tolerances that are applied, as KernelBench uses a uniform threshold across an output tensor, whereas CANNBot applies a dynamic dual-gate precision test [29,78].

Fair Evaluation Challenge #2: Individuals may use a custom harness with different settings rather than the harness released with the benchmark. Even though the benchmark data is the same, settings contained in the harness can have a significant impact on results.

To alleviate this concern and make evaluations comparable, the community should strive to use the harnesses released with the benchmark as much as possible when publishing and comparing quantitative results. If the benchmark’s companion harness is not used, the authors should aim to publiclly release the harness code being utilized. Longer term, the field can strive to provide a standardized evaluation harness that accommodates different hardware platforms. Similar standardization has occurred in other areas, such as object detection in computer vision, where the COCO evaluation protocol (provided in pycocotools) became the standard for evaluating object detection, segmentation, and keypoint detection results, regardless of the imagery employed [79,80].

4.4. Evaluation Integrity Policy

LLMs are creative with respect to the solutions they produce, especially after undergoing RL training. Oftentimes, LLMs will provide a “working” solution to a problem, even if it arrived at that solution through a shortcut. As an example, in Triton kernel generation, if an LLM can only derive solutions that are slower than the PyTorch baseline, without any exploit-detection policies, the model may simply return the PyTorch kernel itself. With this approach, the LLM knows that the kernel is correct and that there is no speed degradation, but this solution is not one that is actually desired. To avoid and penalize such trivial solutions, evaluation pipelines typically include an integrity policy (also informally called “anti-hack” or “anti-cheating” policies), which defines what constitutes as a valid solution versus a hack. This integrity policy should be consistent with the actual prompt that is used to run the LLM, so that the task definition is clear.

With evaluation integrity policies, there is sometimes no standard definition of what is “correct” and “incorrect”, as this may depend on the downstream application or, for corporations, customer requirements. A concrete example can be drawn from adapting the KernelBench benchmark to Triton kernel generation: Is a kernel valid if it still contains numerical operations in PyTorch, or do all numerical operations need to be implemented in Triton? This question poses a subtle, but significant, distinction. Solutions that are required to be pure Triton [32,63] are subjected to a more challenging task — all mathematical operations must be converted to Triton and the overall runtime should be faster than the PyTorch reference. On the other hand, solutions that allow a hybrid mix of Triton and PyTorch [30,54,56] address an easier scenario: If an operation cannot be easily optimized with Triton, then the operation simply remains unchanged as PyTorch. Despite these slightly different problem variants, quantitative results from both are treated equally and compared in the current literature. Figure 12 shows a pictorial example of a hybrid PyTorch + Triton solution on the left and a pure Triton solution on the right.

More generally, evaluation integrity policies extend in several additional dimensions beyond code fallbacks to PyTorch. The four dimensions considered in this survey are summarized in Table 9 and are described here as follows: 1. Fallback checks: As already exemplified above, these checks consider the situation where the model-derived code attempts to use high-level frameworks or libraries as a shortcut in place of a true solution. 2. Generalization checks: Identifies situations where the LLM generates code that is overfit to inputs with specific shapes and values or overfits to the test set. For example, the implementation should not return precomputed constants, or use formulas that are tailored to known tests values rather than computing an operation more generally. 3. Numerical-validity checks: Prevents the model from abusing tolerances or precision during testing. For instance, the model should not produce approximate or low-precision outputs that pass loose tolerance tests, but would be numerically invalid in realistic settings. 4. Measurement-validity checks: Ensures that timing, synchronization, or hidden work manipulations are not performed to make the generated kernel appear more performant. As an example, the model should not attempt to only time a portion of the computation, or move work outside of the measured timed region to improve the generated kernel’s timing profile.

Fair Evaluation Challenge #3: Evaluation integrity policies can significantly affect quantitative results, since they flag solutions as valid or invalid in a binary fashion and even affect the detailed problem definition given to the LLM.

To ensure a fair and comparable evaluation, authors should strive to use the correctness check tools that are distributed with the benchmark. Failing that, when presenting quantitative results, the specific requirements of a correct solution should be defined and the evaluation integrity policy code should be released publicly.

4.5. Evaluation Metrics and Settings

In general, the two main evaluation criteria for automatic kernel generation are: 1) How much faster the output implementation is than a reference, and 2) How many output kernels actually work correctly. As with most fields, no metric perfectly captures all nuances of a problem; as such, a variety of metrics have been proposed to capture the speedup and correctness. The intent of this section is not to highlight every proposed metric; rather, we aim to introduce the dominant metrics while also identifying common pitfalls that can occur, even if the metrics seem clearly defined.

For a pure correctness measure, the most common metric seen in the literature is pass@k. First, consider the case of computing pass@k for a single task (i.e., implementing a single kernel). In this case, the typical unbiased estimator is formally defined as:

$$\widehat{\mathrm{pass}@k}=1-{\displaystyle \frac{\left(\genfrac{}{}{0pt}{}{n-c}{k}\right)}{\left(\genfrac{}{}{0pt}{}{n}{k}\right)}},$$

(2)

where n is the number of candidate kernel generations and c of them are correct. This single-task formulation estimates the probability that if k kernels are drawn from n candidates (with $n\ge k$), at least one of them is correct. However, in most situations, evaluations are not performed over a single task, but across an entire benchmark of kernel tasks. In that case, aggregation across the tasks is required, which can be denoted as:

$$\mathrm{pass}@k={\displaystyle \frac{1}{N}}{\displaystyle \sum _{i=1}^N}{\widehat{\mathrm{pass}@k}}_i.$$

(3)

With respect to measuring the speedup, one of the most basic measurements simply takes the ratio of the baseline execution time divided by the generated kernel’s execution time, which can be expressed as:

$$\mathrm{speedup}={\displaystyle \frac{\mathrm{Baseline}\mathrm{Wall}\text{-}\mathrm{Clock}\mathrm{Time}}{\mathrm{Generated}\mathrm{Kernel}\mathrm{Wall}\text{-}\mathrm{Clock}\mathrm{Time}}}.$$

(4)

Typically, speedup is only computed for those kernels that are designated as being correct. Also, the baseline execution time is usually measured as the reference PyTorch implementation in eager mode, but some papers use torch.compile as the reference time instead. The quantitative results depend directly on what is selected as the baseline implementation being employed.

Building upon simple speedup, KernelBench [29] popularized the fast_p metric, which represents the fraction of tasks where a model produces at least a single generation that is simultaneously correct and achieves a relative speedup threshold of p. The computation for fast_p is as follows:

$${\mathrm{fast}}_p={\displaystyle \frac{1}{N}}{\displaystyle \sum _{i=1}^N}\mathbb{1}\left({\mathrm{correct}}_i\wedge ({\mathrm{speedup}}_i\ge p)\right),$$

(5)

where i iterates over the total number of kernel generation tasks, N. Common values for p are 0.6, 1, 1.2, and 2 (when $p=1$, the metric captures the percentage of tasks where the LLM is able to generate a kernel that matches or exceeds the baseline speed).

The fast_p metric is commonly extended [29,52] to accommodate test-time sampling workflows, whereby kernels are sampled from the model under a consistent and independent procedure. Under such sampling, the model is presented with the same prompt and decoding settings and there is no interaction across the attempts. The rationale for this metric is that in practice, there is often inference-time budget to explore multiple solutions, and sampling can boost the fast_p result. Formally,

$${\mathrm{fast}}_p@k={\displaystyle \frac{1}{N}}{\displaystyle \sum _{i=1}^N}\mathbb{1}\left[\underset{j\in \{1,\dots ,k\}}{\max }\left(\mathbb{1}\left[{\mathrm{correct}}_{ij}\right]·\mathbb{1}[{\mathrm{speedup}}_{ij}\ge p]\right)=1\right].$$

(6)

In words, Equation (6) indicates that for each task, k candidates are considered and if any of them are both correct and sufficiently fast, that task is considered a success. The outer loop over N computes the average across all tasks. Note that a distinction should be drawn here contrasting sampling vs. multi-turn refinement, which is commonly used in Taxonomy Categories #2 - #4. As defined above, sampling produces multiple, independent solutions to a single task. Multi-turn refinement, on the other hand, iteratively invokes an LLM to revise a solution based on feedback from the previous attempts. Multi-turn refinement is a fundamentally different approach and by construction, the turns are not independent.

The metrics pass@k, speedup, and ${\mathrm{fast}}_p$ are arguably the most common used in the field of kernel generation, and most papers do an acceptable job of describing which metric is being deployed. However, there are a number of additional configurations and settings that further confound the task of fair evaluation:

1.

Benchmark Levels. Most benchmarks in the area of kernel generation comprise different levels, roughly corresponding to the size and complexity of the kernel. Papers may report results on a subset of levels and the subsets selected may not be the same among comparable methods.

2.

Baseline timing. The speedups that can be achieved depend directly on the speed of the baseline itself. If the baseline is artificially slow, then large speedups can be achieved, but if the baseline is already fast, further performance increases will be challenging. A recent paper noted that the baseline implementation used by default in KernelBench evaluations was unrealistically slow because TF32 Tensor Core acceleration was disabled [81]. Enabling this setting for the PyTorch baseline saw the overall speedup drop below 1.0 to $0.88\times$ (from $1.43\times$).

3.

Sampling. Whether or not inference-time sampling is employed and if it is, the number of independent samples that are used. A sampling-based approach tends to improve the results [29], so trying to compare two methods with different sampling settings is not fair.

4.

Accelerator hardware and libraries. The specific GPU/NPU/CPU on which the experiments are run, as well as the software libraries, can also affect results. Newer libraries may have implemented improved optimizations, thus making the problem more difficult for an LLM to achieve speed gains. Further, specific hardware configurations may lead to different optimization opportunities. The QiMeng-Kernel paper [40] provided an insightful comparative analysis for Triton kernel generation across three generations of NVIDIA GPUs (V100, A100, and H100). The results showed that when using their method with Gemini 2.5 Pro, fast₁ for Level 1 of KernelBench ranged from 43% on V100 to 67% on H100! Similarly, the mean speedup for the same data was 1.94 on V100, 2.08 on H100, and 2.20 on A100. Thus, for the cleanest comparison, hardware and library settings should all be held constant, but such a requirement is not realistic for researchers across different institutions. Instead, researchers should strive to report all such experimental details and hold as many settings constant as possible.

5.

Last, Average, or Best. As discussed in Section 3, many solutions to kernel generation take an iterative or multi-round approach. Often, the kernel generated during the last turn is the best, but that is not guaranteed. For instance, the LLM may decide to explore a new idea without knowing that it has almost exhausted its turn budget and ends up generating a kernel that does not even compile during its final turn. For this reason, several works consider either the “best” or “average” evaluation settings. Under these protocols, the metrics are computed based on the best kernel that is generated for each task, regardless of its placement in the trajectory or the scores are averaged across all kernels throughout the trajectory.

Fair Evaluation Challenge #4: There are several metrics that are commonly used for kernel generation, however, even if the same top-level metric is used, there are multiple lower-level details and settings that will affect results. To ensure a fair comparison, these details should be reported and matched as much as possible between evaluations.

4.6. Fair Reporting Requirements

Taking the findings of Section 4.1, Section 4.2, Section 4.3, Section 4.4 and Section 4.5 into account, we recommend that papers in the area of kernel generation and optimization report the following experimental details or code to support fair and comparable evaluations.

• Benchmark dataset and any modifications from the original distribution.
• Evaluation harness version (if from a benchmark) or released code.
• Integrity policy version (if from a benchmark) or released code, including numerical threshold tolerances for correctness checks.
• Top-level metrics.
• Benchmark levels (if applicable).
• Baseline reference timing details (e.g., PyTorch eager, torch.compile and whether mixed-precision is permitted for the baseline).
• Generation-time compute budget (e.g., # of LLM calls, total tokens, wall-clock time), since this setting can significantly impact the quality of the generated kernel.
• Number of samples used (assumed 1 by default).
• Exact accelerator hardware (e.g., Ascend 910C, NVIDIA B100).
• Applicable library versions (e.g., CANN, torch_npu, Triton-Ascend, vLLM-Ascend, Mindspore, PyTorch, CUDA toolkit, NVIDIA driver, cuDNN, cuBLAS, etc.).
• Kernel selection scheme, if using a multi-turn or iterative system (e.g., last, average, best, or custom).

These reporting requirements are heavy, but when presented in a compact table, they need not take up excessive space. They could also potentially be reported in an appendix rather than the main paper body. More importantly, inclusion of such experimental details is required to achieve fair and equitable evaluations and monitor which approaches truly advance the state of the art.

Long term, researchers could potentially organize an official online evaluation server, where solutions are independently executed and/or evaluated by an objective third-party. In such a scenario, LLM-generated kernels for a benchmark could be uploaded in a standard format, where timing and evaluation are performed on third-party servers. Results could then be publicly displayed via an online leaderboard. Such an online service would be highly valuable to the community, but it has several drawbacks: 1) It requires significant setup and maintenance efforts. 2) It would require ongoing expense to support timing evaluations from all over the world. 3) It may restrict the acceptable hardware/library configurations that are supported (e.g., it may not consider both NPU and GPU, only cover a specific library configuration, etc.). Due to these challenges, the community’s best path forward in the short term is likely more precise reporting of the experimental setup.

4.7. Comparable Results Summary

Table 10 was compiled to concretely illustrate the challenges facing the field with respect to performing fair evaluations. In this table, a sampling of approaches are included that evaluate on NVIDIA GPUs with the KernelBench benchmark using the same set of first-order metrics. The “Settings” column is included to indicate lower level experimental settings that were used in the following format: (Hardware Platform) / (k, from pass@k) / (one of average, best, last if multi-turn trajectories were used). Note that for some results, rather than report the k value from pass@k, the number of samples was presented instead and is indicated in the table by a _‡ symbol. These two metrics are subtly different, as TritonRL will compute a performance score if any of the 10 samples achieved success, whereas Dr. Kernel’s sampling approach considers the percentage of attempts that were correct. As the “Settings” column demonstrates, there is little consistency across papers with respect to the lower-level evaluation settings. For instance, in multi-turn scenarios, Kernel Foundry [63] only considered the “best” setting, whereas Dr.Kernel primarily used “last”, but did include one result under the “best” configuration. Further, across the eight unique publications in this table, there are at least four different GPUs employed in the experiments. Beyond the information in the “Settings” column, the evaluation harnesses and integrity policy for cheat detection are also not consistent across these works. In fact, several of the papers highlight their robust integrity policies as a novel contribution [30,32,63]. As a result of all these differences, such quantitative comparison tables should be considered as a cue regarding overall trends rather than as a rigorous comparison of which approach is superior.

Despite substantial variation in experimental settings, the final “FRR” column provides encouraging evidence of relatively complete reporting in several recent studies. For four of the eight approaches, TritonRL, Dr.Kernel, DRTriton, and EGG, nine or ten of the 11 items in the Fair Reporting Requirements checklist introduced in Section 4.6 were reported. This assessment considers the papers and supplementary materials, with official code also accepted for the evaluation-harness and integrity-policy requirements. Common reporting gaps concern harness versions or released harness code, integrity policies and numerical tolerances, library versions, and generation-time compute budgets. These findings suggest that adoption of the complete checklist is a realistic objective for future work on automatic kernel generation. More complete reporting would support reproducibility and make comparisons easier to interpret, although it would not by itself ensure comparable evaluation conditions.

With the above limitations in mind, Table 10 supports several cautious observations about the reported results. The Kevin-32B and KernelLLM results reproduced by QiMeng-Kernel [40] exhibit relatively low ${\mathrm{fast}}_1$ rates, particularly at Levels 1 and 3. Several approaches employing iterative refinement, search, or agentic workflows report higher Level 2 ${\mathrm{fast}}_1$ rates than the Category 1 direct generation methods AutoTriton and TritonRL [17,32]. For example, the Dr.Kernel variants report Level 2 ${\mathrm{fast}}_1$ rates of 46.0-80.9%, compared to 12-19% for the TritonRL variants. These differences describe the reported results under their respective evaluation settings; they do not establish the relative performance of the methods under a common protocol.

Among the approaches included in the table, EGG, classified in Category 4, reports the highest correctness and ${\mathrm{fast}}_1$ rates at all three benchmark levels. Kernel Foundry and DRTriton, both classified in Category 3, also report high Level 2 ${\mathrm{fast}}_1$ rates of 90% and 92%, respectively. Kernel Foundry uses DeepSeek-V3 [63], whereas EGG combines GPT-5.1 with expert optimization priors, staged optimization, and specialized agents [21]. The reported results motivate further investigation of search-based and agentic optimization, while controlled comparisons remain necessary to determine the contributions of different algorithmic factors and establish reliable performance rankings.

5. Cross-Platform Kernel Generation: An Ascend NPU Case Study

Throughout this survey, a broad perspective has been adopted that considers research across accelerator types and input/output representations. In this section, the focus is narrowed to the Ascend NPU to examine how automatic kernel generation and optimization differ on this platform compared to NVIDIA GPUs. This section elaborates on Ascend NPU hardware, the unique challenges associated with developing algorithms on Ascend NPU, the potential for portability across platforms, and the current state-of-the-art benchmark results on Ascend NPU versus NVIDIA GPUs.

5.1. Ascend NPU Hardware

Huawei Ascend NPUs are AI accelerators that have been heavily optimized for AI training and inference. In this ecosystem, Ascend is the hardware name, CANN is the core software stack, and AscendC is the low-level operator or kernel programming interface. Figure 13 provides an overview of the complete hardware and software technical stack for Ascend NPUs.

As shown in Figure 13, the Ascend NPU technical stack can be visualized as a layered path from high-level deep learning frameworks, such as PyTorch and MindSpore down to execution on Ascend NPU hardware. Framework adapters, such as “torch_npu” and native Ascend backends, bridge framework-level graphs and tensor operations into the CANN ecosystem. CANN provides the central software environment, including graph compilation, operator libraries, custom-kernel tooling, and debugging support. Although the figure depicts the custom operator/kernel layer and runtime execution as separate stages, these should be interpreted as components that are integrated within the broader CANN ecosystem rather than fully independent layers. The figure is intended as a conceptual overview of the development and execution path, while the precise boundaries between CANN, custom kernel development, and runtime execution are more interconnected in practice.

The basic compute units on Ascend NPUs are called AI Cores. A single Ascend NPU contains multiple AI Cores, each of which consists of:

• Scalar unit: For looping/indexing, address computation, and control flow.
• Vector unit: Fro single instruction multiple data (SIMD) vector computations.
• Cube unit: For accelerated dense matrix multiplication and accumulation (analogous to Tensor Cores on NVIDIA GPUs).

In terms of memory units, Ascend NPUs include the following components:

• Global Memory (GM)/High Bandwidth Memory (HBM): Large off-chip memory (analogous to GPU global memory on NVIDIA GPUs).
• L1 Buffer: On-core local staging buffer.
• Unified Buffer (UB): Local AI Core buffer used frequently for vector operations and intermediate data.
• L0A / L0B: Input buffers closely associated with Cube operations.
• L0C: Output/accumulation buffer for storing the results for Cube computations.
• Registers: For small temporary values, indexing, and scalar control.

Programming on Ascend requires adopting a different mindset compared to NVIDIA GPU programming, focusing on how tensors are moved through specific buffers. For a typical cube computation, the data path is HBM → L1 → L0A/L0B → Cube Unit → L0C → HBM, while for vector computations a standard pathway may be HBM → UB → Vector Unit → UB → HBM.

Figure 14 provides a pictorial overview of how the Ascend 910C NPU compares to an NVIDIA H100 GPU with respect to compute and memory units. At the top architectural level, the AI Core on Ascend NPUs is most comparable to the Streaming Multiprocessor (SM) on NVIDIA GPUs. As mentioned above, Ascend NPUs allocate most computations to either the Cube Unit or Vector Unit, which coarsely correspond to the Tensor Cores and CUDA cores (FP/INT) on GPUs. NVIDIA GPUs typically include warps, which are groups of threads that execute the same instruction together as the basic scheduling unit. A warp scheduler is the hardware inside an SM that selects ready warps to start issuing instructions, which helps hide latency and keep execution units busy. With respect to memory, both platforms have large HBM, but their smaller memories are arranged differently. Ascend NPUs use memory structures that are explicitly segmented into GM, UB, L1, L0A, etc., whereas NVIDIA adopts a cache, shared-memory, and register hierarchy.

The differences in compute and memory architecture between the two accelerator platforms, naturally lead to fundamentally different mindsets and strategies when an engineer or LLM implements kernel operators. On Ascend NPU, the design emphasis is placed on the following questions: 1) How should tensors be tiled? 2) How are tiles distributed across AI Cores? 3) How much data can fit in the UB/L1/L0 buffers? 4) When should data be copied in and results copied out? 5) When should computation be performed? 6) Can data movement and compute be pipelined? In contrast, programming on NVIDIA GPUs focuses more on assigning jobs to threads, maintaining coalesced memory, effectively using shared memory, and ensuring that the Tensor Core is not idling. These differences in hardware and corresponding deviations in design strategy from mainstream NVIDIA GPUs can lead to challenges when developing an automatic kernel generation system for Ascend NPU, as will be discussed next.

5.2. Challenges of Automatic Kernel Generation on Ascend NPU

This section outlines the main challenges of automatic kernel generation on Ascend NPUs. Each subsection states a specific challenge, provides details, and compares it with NVIDIA GPUs, as appropriate.

5.2.1. Challenge #1: Model Generality

As discussed in Section 5.1, for its Ascend NPUs, Huawei has adopted a hardware architecture that is significantly different from NVIDIA GPUs. Accordingly, low-level programming for each platform will require a different set of techniques and skills. As NVIDIA has historically controlled the majority of the hardware accelerator market share in North America, significantly more publications have completed experiments on the GPU platform. SFT and RL checkpoints are commonly shared by researchers following paper publication, making it easier for future contributors to continue investigations along similar directions. For Ascend NPU, fewer such checkpoints are available, and it is also suboptimal to attempt to use a checkpoint fine-tuned with SFT/RL for GPU kernel generation, since the underlying hardware, optimization strategies, and output formats all differ. AscendKernelGen showed that general-purpose, pretrained LLMs such as Qwen3-8B and Llama3.1-8B achieve 0% correctness on NPUKernelBench Level 2 and Level 3, effectively meaning they are unusable without undergoing domain-specific training [33]. A checkpoint fine-tuned for GPU-based kernel generation may be a marginally better starting point than a general pretrained model, but the domain gap between GPU and NPU necessitates fine-tuning to achieve competitive results on Ascend NPU.

5.2.2. Challenge #2: Resource Scarcity

Two critical resources in the deep learning age are data and fine-tuned models. Data is critically important if fine-tuning or post-training is to be performed to adapt to the Ascend NPU problem setting. The number of publicly released kernel generation datasets and benchmarks targeting Ascend NPU is far smaller than that for NVIDIA GPU, as shown in Table 11.

Regarding publicly released checkpoints for models that have undergone SFT or RL for NPU kernel generation, since fewer researchers are working on Ascend devices, the number of such checkpoints is limited. Moreover, as mentioned previously, checkpoints finetuned for GPU kernel generation do not generalize effectively to the NPU platform. To provide concrete evidence of this disparity, at least five different publications cited in this survey have released GPU-oriented checkpoints [17,30,37,38,82], whereas the number of NPU-based checkpoints is far less [33].

Table 11. Public datasets and benchmarks for automatic kernel generation on Ascend NPU versus NVIDIA GPU.

Resource Type	Ascend NPU	NVIDIA GPU
Datasets	Ascend-CoT [33]	KernelBook [83]; CUDA-Agent-Ops-6K [31]
Benchmarks	NPUKernelBench [33]; MultiKernelBench [69]	KernelBench [29]; CUDABench [70]; TritonBench [71]; TritonGym [84]; ComputeEval [85]; FlashInfer-Bench [72]; MultiKernelBench [69]

Table 11. Public datasets and benchmarks for automatic kernel generation on Ascend NPU versus NVIDIA GPU.

Resource Type	Ascend NPU	NVIDIA GPU
Datasets	Ascend-CoT [33]	KernelBook [83]; CUDA-Agent-Ops-6K [31]
Benchmarks	NPUKernelBench [33]; MultiKernelBench [69]	KernelBench [29]; CUDABench [70]; TritonBench [71]; TritonGym [84]; ComputeEval [85]; FlashInfer-Bench [72]; MultiKernelBench [69]

In addition to datasets, benchmarks, and checkpoints, more general resources surrounding the NVIDIA GPU ecosystem are more readily available. For example, there are more sample implementations of GPU kernels across the internet than for Ascend NPU. This situation is similar for toolkits, knowledge bases, and documentation, as well. This lag in ecosystem-level resources puts Ascend NPU automatic kernel development at a technical disadvantage in two ways: 1) Pre-trained models trained on internet-scale knowledge will have better general understanding of NVIDIA GPUs due to the larger amount of content. 2) As agentic workflows become prevalent, the number and maturity of tools/APIs that can be leveraged by agents is greater on the NVIDIA platform.

This imbalance in resources current presents an additional challenge for automatic kernel generation on Ascend NPU. However, it is expected that as the Ascend NPU user base of researchers and industrial partners grows, the ecosystem will also mature, leading to a shrinking gap compared to NVIDIA GPUs.

5.2.3. Challenge #3: Output Formats

The introduction of higher-level DSLs, such as Triton-Ascend, has made many aspects of kernel development easier and more accessible compared to lower-level interfaces such as AscendC and CUDA. Ascend NPUs and NVIDIA GPUs support Triton-Ascend and Triton, respectively, which share significant syntactic similarities. As a result of these similarities, generating a working kernel on either platform is typically a task of comparable complexity. However, the same cannot be said for performance (i.e., speed improvements) on NPU versus GPU. Typically, speedup results on NPU have been lower than those achieved on GPU [69]. One reason for this discrepancy is that Triton was originally shaped around NVIDIA-style GPU programming, while Triton-Ascend has mapped that programming model onto the Ascend NPU hardware and CANN compiler stack. In other words, Triton was designed for GPUs, and Triton-Ascend was adapted to the Ascend NPU platform, making it a less natural fit.

When considering beyond high-level DSLs to lower-level AscendC and CUDA, the syntactic differences become more pronounced, as shown in Figure 15. Specifically, AscendC consists of multiple files, both on the host CPU side and on the NPU device side. These files include a tiling function and a host launcher, as well as a header file and the actual kernel implementation on the device. In contrast, a typical CUDA kernel often requires just a single, self-contained .cu output file. Creating a complete operator package for AscendC is more challenging for existing LLMs, since an understanding is required of the various files and their respective roles [69]. This challenge is exacerbated by the data and resource scarcity issue discussed in Challenge #2—without a sufficient number of training examples, an LLM is not able to learn and understand the roles and interrelationships between the different host and device files. In principle, modern LLMs are capable of producing multi-file artifacts and thus, as additional training resources become available, the gap in performance between AscendC and CUDA kernel generation is likely to reduce.

5.3. NPU Kernel Generation Performance

To exemplify the aforementioned challenges, this section provides specific references to state-of-the-art performance improvements on Ascend NPU compared with NVIDIA GPU, considering both Triton-Ascend and AscendC.

5.3.1. Triton

For Triton code generation, AKG Kernel Agent is the only work to our knowledge that performs experiments directly comparing experimental results with both Ascend and CUDA backends [66]. Table 12 shows the results originally presented in Du et al. [66] for kernel correctness with Triton-Ascend compared to Triton, where DeepSeek V3.1 (non-reasoning mode) was employed. An Ascend 910B NPU was used while the GPU was an NVIDIA A100. As the table shows, with respect to correctness for Triton on GPU, the model can generate working code for all operator categories, while the LLM struggles to even generate working code for some operator types on NPU. The convolution category contains the most examples (34 operators) and has the lowest correctness score of 41.2% on NPU.

For performance, Table 13 shows the geometric mean speedup and fast@k metrics with the same experimental setup as discussed above from Du et al. [66]. Considering only the final “Overall” row of this table can be misleading, as it suggests that the geometric mean speedup is notably higher with Triton-Ascend than Triton. However, closer inspection reveals that this higher result was achieved because the Convolution category was evaluated for Triton but omitted for Ascend NPU. As the GPU-based Triton numbers illustrate, Convolution is the most challenging operator type and the performance numbers are significantly lower for that category than any other on GPU. In fact, in a side-by-side comparison, the only result that is higher for Triton-Ascend is the geometric mean for the “Reduce & Norm” category (1.66 on NPU vs. 1.44 on GPU).

These results suggest that, under the experimental conditions reported in AKG Kernel Agent [66], LLM-generated Triton-Ascend kernels were less consistently successful than their Triton-CUDA counterparts. However, this comparison should not be interpreted as isolating the effect of the programming model alone, since the results inherently include differences in hardware, compiler maturity, runtime stack, benchmark design, and backend support. Disentangling these factors remains an open problem for future cross-platform kernel-generation benchmarks.

5.3.2. AscendC

Several recent studies have demonstrated that a significant quantitative gap in performance exists when applying state-of-the-art foundation models to the task of kernel generation in AscendC compared to CUDA [39,69,86]. For instance, Table 14 shows results originally reported in Wen et al. [69] across seven high-profile pre-trained models on MultiKernelBench. In this paper, the authors introduce metrics that are closely related to pass@k and fast_p. Specifically, the authors utilize the following metrics: 1) Compilation@k, which measures the percentage of tasks where one or more of the k generated kernels compiles successfully. 2) SpeedUp_α@k (SU_α@k), which denotes the percentage of tasks where one or more samples reach a speedup of $\alpha$ or greater and 3) Pass@k, which indicates the percentage of tasks where one or more samples produce the correct, expected output. Automatic CUDA kernel generation enjoyed successful compilation rates as high as 97.5% with GPT-4o, whereas the highest compilation rate for AscendC was only 10.2% with DeepSeek-V3-0324. Moreover, DeepSeek-V3-0324 achieved a SU₁ for just over 1% of the AscendC kernels, whereas the top-ranking models obtained a SU₁ for 26% of the kernels when using CUDA.

Results from Zheng et al. [39] report similar outcomes, as shown in Table 15. With GPT-5.2, for example, the number of correct CUDA kernels is 78% points higher than AscendC on Level 1 (92% vs. 14.0%) and 88% points higher on Level 2 kernels (90% vs. 2%). The current gap with state-of-the-art pre-trained foundation models on Ascned NPU compared to NVIDIA GPU is enormous.

These studies have shown that despite many of the state-of-the-art pre-trained models achieving success on CUDA kernel generation, they remain largely unusable and generate primarily incorrect kernels for AscendC. These results speak to the challenge of resource scarcity, in that large foundation models were not sufficiently exposed to the task of AscendC kernel development in their internet-scale pre-training.

6. Future Directions and Opportunities in Automatic Kernel Generation

Throughout this survey, several challenges have been highlighted with respect to automatic kernel generation. These challenges can be viewed as future directions or open research opportunities, as will be discussed next.

6.1. Fair Evaluation

A theme expressed throughout this survey, especially in Section 4, is that it is challenging to perform fair comparisons across the current set of research publications to conclude which methods perform best. Several factors make fair evaluations complicated, including different hardware platforms, benchmarks, metrics, and evaluation harnesses. In this survey, we have proposed reporting guidelines to assist the field in organizing published quantitative results and making them more directly comparable. It is now up to researchers in the field to adopt and refine such guidelines as needed.

An even greater opportunity for fair evaluation is for a team of researchers to create an objective, third-party evaluation platform with an online leaderboard, similar to what has been done in other fields, such as robotics [87]. Setting up such an evaluation platform is costly and requires significant effort, but it would be a major contribution to the field and promote fair, objective evaluations.

6.2. Improved Performance on Higher-Level Tasks

Current methods typically exhibit significantly stronger results on lower benchmark levels that focus on isolated operators. In comparison, higher benchmark levels (e.g., KernelBench Level 3) often consider larger fused modules or model fragments, and performance on such scenarios has so far been relatively poor. For higher-level tasks, an LLM must be able to rewrite a non-trivial part of a model architecture without corrupting the computation, while also identifying efficiencies beyond an already optimized framework. That task is arguably much harder than generating and optimizing isolated kernels. There is still significant room for methods to advance the state of the art for these higher levels.

6.3. Training-Based Approaches Merged with Advanced Pipelines

This survey has presented a taxonomy that separates methods based on their generation-time workflows in order to disentangle approaches based on whether they are “training-based”, “training-free”, or “agentic”. The taxonomy focuses on inference-time harnesses, while presenting whether an approach leverages any training techniques as an orthogonal axis. From the taxonomy, one observation is that there is a lack of approaches that leverage specialized, kernel-generation-focused training within Category #4 of the taxonomy (adaptive tool-orchestrating agents). This finding is shown clearly in Table 7. Even within Category #3, search-based kernel generation, only three approaches [54,58,64] have considered any type of learning in their pipelines. In either case, there are still significant opportunities to explore fine-tuning of the LLM for specialized kernel-generation tasks within an agent-based or search-based framework that is employed during generation.

6.4. Industry-Oriented Open Problems

6.4.1. Refined Metrics, Objectives, and Reward Functions

As shown in this survey, evaluation metrics mostly target correctness and speedup, and therefore the losses used when fine-tuning models typically focus exclusively on these two factors. In industrial practice, there are a multitude of additional dimensions that should be considered when judging the quality of a kernel implementation, including latency, memory usage, numerical stability, portability, and maintainability. There is significant room to explore metrics and training rewards that consider such additional, industry-oriented factors.

6.4.2. Realistic, Industry-Level Correctness Evaluations

Most evaluation harnesses include correctness checks to ensure that the kernel actually produces the desired output before subjecting it to more rigorous profiling. However, for some benchmarks (e.g., [29]), the correctness check is rudimentary, considering only a single set of input values with a specific shape. Kernel operators are effectively functions, and it is not realistic to conclude that a function has been implemented properly from a single data point. In Cao et al. [33], a more complete correctness check is performed that includes an average of 50 inputs with different values, but the shapes of the inputs remain unchanged throughout. KernelBench-Verified [81] also recently identified a limitation of the standard KernelBench evaluation [29], whereby models were learning implementation shortcuts based on the distribution of input values used during evaluation. As a solution, KernelBench-Verified proposed using multiple distributions of input values, including larger, smaller, and negative values. In an industrial setting, these types of issues become operational risks because the generated kernels may be executed automatically inside larger systems.

NPUKernelBench [33] and KernelBench-Verified [81] are steps in the right direction, but current and future benchmarks should be expanded to consider different data types and edge cases as well as define production-grade evaluation integrity policies. Moving forward, it is possible that benchmarks and solutions may diverge into two sets: Those that consider static-shape kernel specialization and those that consider dynamic-shape or robust operator generation. Most existing benchmarks evaluate static-shape specialization [29,69,70,71], but some benchmarks already consider dynamic-shape evaluations [14,66]. Benchmarks that specifically target a method’s ability to generate kernels that accommodate dynamic shapes is also certainly an area for future expansion.

6.4.3. Production Integration

The industrial value of a kernel implementation is ultimately dictated by whether it can be integrated into a real inference or training stack. A generated kernel must adhere to framework APIs, build systems, dispatch logic, shape handling, and code-review expectations. Kernel-Smith is the only work at this time that explicitly studies whether generated kernels can be turned into contributions on SGLang and LMDeploy [64]. This seminal work suggests that future industrial-focused research should not only evaluate benchmark speedup, but also whether generated kernels can be accepted into production codebases.

6.4.4. Technical Stack and Version Generalization

Industrial kernel performance is highly dependent on the underlying hardware, compiler, driver, runtime, and framework stack. Some Ascend NPU-oriented papers have reported results with precisely-specified stacks [14,35]. AscendOptimizer reports that an Ascend 910B2 with CANN 8.3 were utilized [14], while AscendCraft explicitly denotes the NPU, CANN, PyTorch framework, and OS versions [35]. For an automatic kernel generation system to be used in practice, an industrial user needs confidence that a generated kernel will remain correct and performant after toolchain upgrades, firmware changes, or deployment on similar accelerator variants. Key open problems are therefore robust kernel generation across platform variants and evaluating such performance, not merely achieving peak performance on a singular experimental setup.

6.5. Ascend NPU-Specific Open Problems

6.5.1. Cross-Platform Generalization

Current evidence from state-of-the-art pretrained models suggests that although they can achieve reasonable performance for CUDA kernel generation, they experience a collapse in performance when generating AscendC [39,69,86]. This result suggests that models cannot currently generalize across hardware platforms. This open problem may be longer-term in nature, since the field is still focused on improving performance for individual hardware platforms; however, the development of a fully generalized, cross-platform automatic kernel generation system presents a significant opportunity with a profound impact on the field.

6.5.2. Lagging Performance on Ascend NPU

Another point reiterated throughout this survey is that results on Ascend NPU tend to lag behind those attained on other platforms, both at the high-level DSL level (e.g., Triton and Triton-Ascend) and the lower level (e.g., AscendC and CUDA). This performance gap requires additional analysis to determine whether it is purely a result of limited datasets, benchmarks, and other resources, or whether there are more fundamental causes. For example, at the Triton-Ascend level, is part of the performance gap a result of the fact that Triton was first designed and customized for CUDA hardware and only later adapted to Ascend NPU? Does this situation limit the performance ceiling that can be achieved with Triton-Ascend? Similarly, is AscendC’s multi-file operator package fundamentally a more challenging problem for an LLM to solve than producing a single CUDA output file?

In addition to the above analyses, new search or optimization solutions may be required for AscendC generation due to the need to co-optimize both host-side and device-side implementations. Exploration of such co-optimization methods is another potential direction for future work.

7. Summary and Conclusions

Building on the significant success that LLMs have shown for general code generation and debugging, this survey has reviewed the body of work on automatic kernel generation and optimization. Similar to the general coding scenario, automatically generating kernel code offers significant advantages in terms of development time and cost. The benefits are even more pronounced for kernel generation, since it is a highly specialized, complex task that requires engineers that possess a deep understanding of both software and the underlying hardware platform. Accelerating this development process through partial or full automation could lead to significant competitive advantages for companies.

As a result of the significant benefits of automatic kernel generation, the field has experienced rapid growth in the past two years, as outlined in this survey. Due to this rapid development, the field has encountered challenges, including a lack of organization and equitable evaluations, but the potential of automatic kernel generation remains clear. In this survey, we provide much-needed organization by presenting a novel taxonomy that focuses on the generation-time harness workflow and present a detailed analysis of the challenges that prevent researchers from producing fair, comparable evaluations. To address the evaluation challenges, we propose solutions and provide a list of standard items that should be reported when publishing results. This survey also considers multi-platform generalization and provides a detailed case study on the Ascend NPU platform, considering its hardware, programming languages, domain-specific languages, and how these aspects affect observable performance trends. Despite the rapid progress, many open problems remain in the field, and these issues were discussed in Section 6.

We hope that this survey provides a useful snapshot for researchers in the field, highlighting the current state of research, the remaining goals, and possible paths toward achieving those goals. The future of industrial kernel design and optimization will likely involve a synergistic relationship between human experts and AI-powered tools, leading to more efficient and shorter development cycles, as well as lower training and inference costs for organizations and end users.