Paper deep dive

Towards Cold-Start Drafting and Continual Refining: A Value-Driven Memory Approach with Application to NPU Kernel Synthesis

Yujie Zheng, Zhuo Li, Shengtao Zhang, Hanjing Wang, Junjie Sheng, Jiaqian Wang, Junchi Yan, Weinan Zhang, Ying Wen, Bo Tang, Muning Wen

Year: 2026Venue: arXiv preprintArea: cs.LGType: PreprintEmbeddings: 95

Intelligence

Status: succeeded | Model: google/gemini-3.1-flash-lite-preview | Prompt: intel-v1 | Confidence: 95%

Last extracted: 3/22/2026, 6:18:35 AM

Summary

EvoKernel is a self-evolving agentic framework designed to automate kernel synthesis for data-scarce Domain-Specific Architectures (DSAs) like NPUs. It addresses the 'cold-start' problem by formulating kernel synthesis as a memory-based reinforcement learning task, utilizing a value-driven retrieval mechanism to learn stage-specific Q-values for drafting and refining kernels without requiring expensive fine-tuning.

Entities (5)

EvoKernel · framework · 100%Ascend C · programming-language · 95%KernelBench · benchmark · 95%Large Language Models · technology · 95%NPU · hardware-architecture · 95%

Relation Signals (3)

EvoKernel → automates → kernel synthesis

confidence 95% · EvoKernel, a self-evolving agentic framework that automates the lifecycle of kernel synthesis

EvoKernel → evaluatedon → KernelBench

confidence 95% · By building an NPU variant of KernelBench and evaluating on it, EvoKernel improves frontier models' correctness

EvoKernel → utilizes → value-driven retrieval

confidence 95% · Through a novel value-driven retrieval mechanism, it learns stage-specific Q-values

Cypher Suggestions (0)

No Cypher suggestions yet.

Abstract

Abstract:Deploying Large Language Models to data-scarce programming domains poses significant challenges, particularly for kernel synthesis on emerging Domain-Specific Architectures where a "Data Wall" limits available training data. While models excel on data-rich platforms like CUDA, they suffer catastrophic performance drops on data-scarce ecosystems such as NPU programming. To overcome this cold-start barrier without expensive fine-tuning, we introduce EvoKernel, a self-evolving agentic framework that automates the lifecycle of kernel synthesis from initial drafting to continual refining. EvoKernel addresses this by formulating the synthesis process as a memory-based reinforcement learning task. Through a novel value-driven retrieval mechanism, it learns stage-specific Q-values that prioritize experiences based on their contribution to the current objective, whether bootstrapping a feasible draft or iteratively refining latency. Furthermore, by enabling cross-task memory sharing, the agent generalizes insights from simple to complex operators. By building an NPU variant of KernelBench and evaluating on it, EvoKernel improves frontier models' correctness from 11.0% to 83.0% and achieves a median speedup of 3.60x over initial drafts through iterative refinement. This demonstrates that value-guided experience accumulation allows general-purpose models to master the kernel synthesis task on niche hardware ecosystems. Our official page is available at this https URL.

PDF

Open source PDF →Open local PDF →

Trouble viewing inline? Open PDF directly →

Full Text

94,731 characters extracted from source content.

Expand or collapse full text

TOWARDS COLD-START DRAFTING AND CONTINUAL REFINING: A VALUE-DRIVEN MEMORY APPROACH WITH APPLICATION TO NPU KERNEL SYNTHESIS Yujie Zheng *, 1 , Zhuo Li *, 1 , Shengtao Zhang 1 , Hanjing Wang 2 , Junjie Sheng 3 , Jiaqian Wang 1 , Junchi Yan 1 , Weinan Zhang 1 , Ying Wen 1 , Bo Tang 4 , Muning Wen †, 1 1 Shanghai Jiao Tong University 2 Shanghai Artificial Intelligence Laboratory 3 Independent Researcher 4 MemTensor (Shanghai) Technology Co., Ltd ABSTRACT Deploying Large Language Models to data-scarce programming domains poses significant challenges, particularly for kernel synthesis on emerging Domain-Specific Architectures where a “Data Wall” limits available training data. While models excel on data-rich platforms like CUDA, they suffer catastrophic performance drops on data-scarce ecosystems such as NPU programming. To overcome this cold-start barrier without expensive fine-tuning, we introduce EvoKernel, a self-evolving agentic framework that automates the lifecycle of kernel synthesis from initial drafting to continual refining. EvoKernel addresses this by formulating the syn- thesis process as a memory-based reinforcement learning task. Through a novel value-driven retrieval mechanism, it learns stage-specific Q-values that prioritize experiences based on their contribution to the current objective—whether bootstrapping a feasible draft or iteratively refining latency. Furthermore, by enabling cross-task memory sharing, the agent generalizes insights from simple to complex operators. By building an NPU variant of KernelBench and evaluating on it, EvoKernel improves frontier models’ correctness from 11.0% to 83.0% and achieves a median speedup of 3.60×over initial drafts through iterative refinement. This demonstrates that value-guided experience accumulation allows general-purpose models to master the kernel synthesis task on niche hardware ecosystems. Our official page is available at https://evokernel.zhuo.li. 1INTRODUCTION A practical limitation when deploying Large Language Models (LLMs) to niche domains is their inability to generalize beyond their pre-training distribution (Minaee et al., 2024; Wang et al., 2025). When faced with cold-start scenarios, domains where training data is sparse and expert demonstrations are unavailable, even frontier models struggle significantly (Kostikova et al., 2025; Joel et al., 2024). This challenge is particularly acute in domains where (i) correctness is binary and machine-verifiable, leaving little room for “partially correct” solutions (Jain et al., 2024; Yan et al., 2024), (i) expert knowledge is scarce and expensive to acquire, and (i) the gap between in-distribution and out-of-distribution performance is stark. Automated kernel synthesis for emerging hardware accelerators exemplifies this extreme scarcity (Yu et al., 2026). While the industry is aggressively diversifying toward Domain-Specific Architectures (DSAs) like NPUs, TPUs, and neuromorphic chips (Silvano et al., 2025; Liao et al., 2021; Jouppi et al., 2023) to address escalating computational costs (Kaplan et al., 2020), these nascent ecosystems face a severe “Data Wall”. Unlike the mature NVIDIA landscape, where decades of CUDA repositories provide a massive pre-training corpus, emerging platforms are characterized by extreme data scarcity: public code is rare, documentation is esoteric, and compiler * Equal contribution. † Corresponding author: Muning Wen (muningwen@sjtu.edu.cn) 1 arXiv:2603.10846v1 [cs.LG] 11 Mar 2026 Table 1: Few-shot functional correctness (pass@4) of frontier LLMs on CUDA vs. Ascend C kernel generation. Results are from our experiments; the level definitions (L1, L2) and setup details are consistent with Section 4.1. ModelLevelCUDA (%)Ascend C (%) GPT-5.2 L192.014.0 L290.02.0 DeepSeek-V3.2 L150.08.0 L29.00.0 Qwen3-Coder-30B L146.07.0 L210.00.0 feedback is opaque (Joel et al., 2024). This barrier is compounded by the fact that highly optimized CUDA kernels (Choquette et al., 2021; Wu, 2023) are not portable to these architectures due to fundamental differences in memory hierarchy and instruction sets, leaving foundation models with virtually no expert demonstrations to bridge the cold-start gap. As evidenced in Table 1, state-of-the-art LLMs that achieve high performance on CUDA (Ouyang et al., 2025) suffer a catastrophic collapse when transferred to a data-scarce Domain-Specific Language (DSL) like Ascend C, which is specifically designed for NPU kernel programming. In line with prior findings (Wen et al., 2025), even GPT-5.2, which attains 92% on CUDA L1 tasks, drops to 14% on Ascend C; on the more challenging L2 tasks, models fail entirely. This observation suggests that current models do not genuinely “learn” to program new hardware like NPUs, but instead rely on memorized patterns from pre-training distributions. Standard paradigms to bridge this gap prove insufficient in such data-scarce domains. Supervised Fine-Tuning (SFT) (Zhou et al., 2023; Chung et al., 2024) demands thousands of expert-labeled examples per domain (Longpre et al., 2023), which is prohibitively expensive when targeting rapidly evolving or niche environments like NPU programming. Parametric policy-based Reinforcement Learning (Zhang et al., 2025; Kakade, 2003) requires extensive online rollouts to update model weights, incurring high sample complexity (Cao et al., 2024; Qi et al., 2025) and risking catastrophic forgetting of general capabilities. Traditional Retrieval-Augmented Generation (RAG) (Lewis et al., 2020) falters when the database is sparse (Contal & McGoldrick, 2024; Barnett et al., 2024); even with relevant samples, similarity-based retrieval does not guarantee effectiveness (Izacard et al., 2023). Consequently, the core challenge is a cold-start problem: How can an agent autonomously master a rigorous, data-scarce kernel synthesis task from scratch, without expert demonstrations or expensive fine-tuning? To address this, we introduce EvoKernel, a framework that formulates kernel synthesis as a reinforcement learning task over a self-evolving memory. By employing a novel value-driven retrieval mechanism, the agent learns stage-specific Q-values to quantify the utility of historical experiences, dynamically shifting focus from bootstrapping functional correctness (Drafting) to optimizing latency (Refining) without updating model weights. Empirically, EvoKernel bridges the cold-start gap on NPU benchmarks, boosting the correctness of frontier models from 11.0% to 83.0% and achieving a 3.60x median speedup over the first feasible draft, thereby demonstrating that value-guided experience accumulation enables general-purpose models to master data-scarce hardware ecosystems. Our contributions are summarized as follows: •Unified Drafting-Refining Pipeline: We propose a two-stage framework over a shared memory that transitions from feasibility-driven drafting to latency-driven refining to bootstrap and optimize NPU kernels. • Evolving Value-Driven Retrieval: We introduce a retrieval mechanism that learns stage-specific Q-values to quantify memory utility. A unified Monte-Carlo update adapts the policy from verifier feedback without updating model weights. • Comprehensive Evaluation and Insights: EvoKernel boosts performance on NPU benchmarks from 11.0% to 83.0%. We provide in-depth analysis of cross-task transfer, emergent curricula, and scaling to out-of- 2 distribution workloads such as the Attention Set and recent MHC kernels, demonstrating how memory autonomously bridges the data-scarce gap. 2RELATED WORK Self-Evolving and Adaptive Agents. While Large Language Models (LLMs) are typically static, recent research explores mechanisms for self-improvement. Inference-time techniques, such as Self-Refine (Madaan et al., 2023) and Tree-of-Thoughts (Yao et al., 2023), utilize iterative critique loops to enhance reasoning within a single episode, though these improvements are transient, resetting once the context window closes (Shinn et al., 2023). Closest to our work are evolutionary frameworks like AlphaEvolve (Novikov et al., 2025) and EvolveR (Wu et al., 2025), which accumulate experience across episodes. These methods typically assume sufficient initial competency or verifiable intermediate states, conditions absent in the rigid “all-or-nothing” compilation environment of data-scarce kernel synthesis, where our approach operates. Memory-Augmented Generation. To overcome context limitations, systems like MemGPT (Packer et al., 2023) and MemOS (Li et al., 2025c; Chhikara et al., 2025) introduce operating-system-like memory hierarchies for long-horizon tasks. In agentic workflows, Voyager (Wang et al., 2023) and other generative agents (Park et al., 2023; Fang et al., 2025) demonstrate the power of retrieving procedural skills or behavioral reflections (Madaan et al., 2022). More recently, Memento (Zhou et al., 2025) and MemRL (Zhang et al., 2026) have formalized retrieval as a reinforcement learning problem, learning what to retrieve. We adapt this value-based retrieval paradigm to kernel engineering, where surface-level semantic similarity often fails. Automated Kernel Synthesis. Kernel synthesis demands strict functional correctness and hardware-specific optimization. Benchmarks like KernelBench (Ouyang et al., 2025) and MultiKernelBench (Wen et al., 2025) reveal that general-purpose LLMs degrade sharply on unfamiliar backends due to domain shifts (Li et al., 2025a). To mitigate this, recent agentic frameworks such as QiMeng-Kernel (Zhu et al., 2025) and KernelBand (Ran et al., 2025) utilize iterative execution feedback for refinement, as do multi-agent systems like STARK (Dong et al., 2025) and AKG Kernel Agent (Du et al., 2025). Supervised approaches like Kevin (Baronio et al., 2025) and AutoTriton (Li et al., 2025b; Woo et al., 2025) fine-tune models on domain-specific corpora. These methods often assume access to high-quality training data, limiting their applicability in emerging ecosystems. EvoKernel addresses this cold-start setting by learning to retrieve from a self-evolving memory bank rather than relying on static corpora. 3EVOKERNEL: VALUE-DRIVEN MEMORY UPDATE FOR KERNEL EVOLUTION As shown in Figure 1, we propose the EvoKernel, a framework that automates the lifecycle of hardware-specific kernel synthesis, from cold-start drafting to continual performance refinement. In this paper, we instantiate the framework primarily on Ascend C, while the same agent loop can be specialized to other backends through backend-specific prompts, verifier toolchains, and profiling signals. We formulate this process as a Memory-based Markov Decision Process (M-MDP) (Zhou et al., 2025; Zhang et al., 2026), where an agent learns to retrieve high-utility experiences to guide a LLM generator. 3.1PROBLEM FORMULATION A kernel synthesis taskx∈Xis specified by a PyTorch reference operator and metadata (e.g., input shapes and operator hyperparameters). Given a taskxand retrieved contextc, a generatorG θ samples a kernel and the goal is to generate a kernel source code y ∈Y that satisfies functional correctness and minimizes execution latency. We model the generation process as an M-MDP over a horizonT.A trajectory is defined asτ = (s 0 ,c 0 ,a 0 ,r 0 ,...,s T ), governed by the tuple (S,A,M,P,R). The components are defined as follows: State Space (S): A states t is defined as a tuple(x,ξ t ), wherex ∈ Xdenotes the static kernel task (PyTorch operator + metadata), andξ t represents the dynamic generation state (e.g., current best-so-far latency or verification status). 3 From Generation Traces, Historical Attempts on the Same Task 푥: Feasible AttemptFailed Attempt 0.671.03-0.04 ... -0.04 Refined Code Continual Refining Cold-Start Drafting Drafted Code Latencyvs. Best -1 Anti-HackingCompilationCorrectness Performance Verifier AscendC APITemplates ValueDrivenMemory Success & Failure Experiences Generation traces Best Practices ComposetoBuild Context APIs Experiences from the Same Task or Other Tasks Contextc t Environment 퓟(풙) ComposetoBuild Context Best practices Attempts from Other Tasks Experiences Retrieved by Q Task Batch 퓧 1.03 퓅 Top-k Retrieved by Q AttemptsfromOtherTasks +1 Value Iteration: 푸=푸+휶풓−푸 Q Q Q Q Retrieved items ·Code · Verified Results · Rewards Child nodes are observable Reward Figure 1: The EvoKernel framework. (Left) Cold-Start Drafting: Given task batchX, retrieves top-kcandidates, filters context viaQ, and synthesizes an initial kernel. (Center) Environment & Memory: A multi-gate verifier assesses generated code to yield rewards, which updateQvia value iteration; code and results are stored in Memory. (Right) Continual Refining: Exploits generation tracesP(x)and historical attempts, including observable child nodes, to iteratively optimize for lower latency. Action Space (A): The action a t ∈A corresponds to a generated kernel code y ∈Y . Memory (M): We defineM t as a dynamic, self-evolving memory bank. It is initialized asM 0 comprising seed knowledge. At each step t, it accumulates the agent’s interaction history, updating according to the rule: M t+1 ←M t ∪(s t ,a t ,r t ),(1) Transition Dynamics (P): The transition dynamicsP :S×A→ ∆(S)describe the evolution of the generation process. Since task x remains invariant within an episode,P deterministically updates the generation state: s t+1 = (x,ξ t+1 ), ξ t+1 = f (x,ξ t ,a t ,o t ),(2) Here,fupdates the dynamic generation state by integrating the actiona t and its verifier outcomeo t , conditioned on the task x and the previous state ξ t . Reward Function (R): The environment provides a scalar feedback signalr t ∈ Rbased on evaluation of the action a t . Policy Factorization. To tackle this M-MDP, the agent operates via a composite policy. At each stept, a Retrieval Policyμfirst selects a contextc t ⊂M t based on the current state. Conditioned on this context, the Generator Policy G θ samples the code: π(y t |s t ,M t ) = G θ (a t |s t ,c t )· μ(c t |s t ,M t ),(3) Our core methodology focuses on optimizingμvia reinforcement learning to identify high-utility memory items, while G θ leverages the pre-trained capabilities of the LLM. 3.2MEMORY ARCHITECTURE AND VALUE-DRIVEN RETRIEVAL The efficacy of the generatorG θ depends critically on the quality of the contextc t . We designMas a heteroge- neous knowledge base containing: (i) API templates for the active backend when such documentation is available (e.g., Ascend C), (i) summarized success and failure experiences, (i) generation traces, including both draft and refined variants, and (iv) best practices for kernel refinement. To instantiate the policyμ, we introduce Value-Driven Retrieval. Unlike traditional similarity-based retrieval, our approach dynamically evaluates memory item utility based on the current generation stage. For statesand candidate memory itemm, we define a Q-value functionQ k (s,m)that estimates the expected benefit of including m in the context at stage k. 4 For a given taskx, letNdenote the final retrieval count. We first use dense retrieval to obtain a top-Kcandidate poolC(x)⊂M, whereK = λNandλis an over-retrieval multiplier. We then use stage-specific value estimates Q k to filter these top-Kcandidates down to the finalNitems for the context. These values reflect the agent’s evolving objectives: • Drafting Stage (Q 1 ): Estimates the likelihood that m contributes to a functionally correct kernel. •Refining Stage (Q 2 ): Estimates the contribution of a memory itemmto latency optimization of the kernel, where m can either be an optimization start point p or auxiliary refinement items fromM. In the upgraded system, the drafting context is assembled by a hybrid retrieval policy: experiential memories and code traces remain value-selected, while API knowledge is retrieved through a backend-aware mixture of static infrastructure bundles, exact-name lookup from retrieved code examples, and semantic/category-based search. This separation is important in practice because API utility is largely determined by backend coverage and operator decomposition, whereas experiential memories benefit directly from online value estimation. Unified Value Update. Despite the distinct objectives, we employ a unified Monte-Carlo (MC) update rule to refine the retrieval policyμ. Upon observing a rewardr t (defined in subsequent sections) after using context items c t , we update the Q-values for all m∈ c t : Q(s,m)← Q(s,m) + α· r− Q(s,m) ,(4) whereαis the step size. This update rule allows the retrieval policyμto continuously adapt to the evolving capabilities ofG θ . We provide formal guarantees on boundedness and convergence of these value estimates in Appendix A. 3.3STAGE 1: COLD-START DRAFTING The objective of this stage is to obtain an initial feasible kernel that can bootstrap subsequent refinement. For a taskx, we iteratively (i) retrieve a drafting contextc t ⊂C(x)using anε-greedy policy overQ 1 , and (i) sample a candidate kernel y t ∼ G θ (·| x,c t ). Reward and update. We use a binary feasibility reward r 1,t = ( +1, if g feas (o t ) = 1, −1, otherwise, (5) whereo t = V (x,y t )andg feas is the combined feasibility gate (Section 3.5). After receiving feedback, we update the values of retrieved entriesm∈ c t using Eq. 4 withr = r 1,t and store the generated code together with verifier feedback into memory. This process repeats until a feasible kernel is found or the budget is exhausted. 3.4STAGE 2: CONTINUAL REFINING Once a feasible kernel is obtained, the focus shifts from feasibility to latency reduction. We maintain a set of optimization start pointsP(x), initialized with the successful draft from Stage 1 and augmented online as new feasible variants are discovered. At each iteration, based on the current states t , we retrieve the available start points from the memoryM and select a start point using Q 2 . With the selected start point and the current state, we then retrieve additional contextual information that contains optimization traces, best practices, and information about its observable child nodes to support the refinement process. In the upgraded system, this refinement context is further conditioned on profiler-derived bottleneck diagnoses, which are used to retrieve bottleneck-matched optimization examples and complementary high- performing variants. Using the selected start point and the retrieved context inc t , the generator samples a refined result. Relative reward, normalization, and update. To drive performance optimization, we define reward relative to the best-so-far latency b t tracked in ξ t : r 2,t = ( −1,if g feas (o t ) = 0, tanh(logb t − logℓ lat (o t )), otherwise. (6) 5 Table 2: Compilation Rate (CR) and Correctness (Acc) across difficulty levels, shown as (Round 1) Final, respectively representing the start point and the final performance. Notably, the huge gap between GPT-5.2 and other models implies that frontier LLMs with stronger in-context learning capability benefit substantially more from experience-driven methods. ModelMethod Level 1Level 2Overall CR (%)Acc (%)CR (%)Acc (%)CR (%)Acc (%) Qwen3-Coder-30B Pass@k(22.0) 30.0(7.0) 8.0(0.0) 2.0(0.0) 0.0 (11.0) 16.0(3.5) 4.0 Refinement (13.0) 22.0(2.0) 6.0(0.0) 1.0(0.0) 0.0(6.5) 11.5(1.0) 3.0 Ours(25.0) 33.0(6.0) 11.0(1.0) 3.0(0.0) 0.0 (13.0) 18.0(3.0) 5.5 DeepSeek-V3.2 Pass@k(21.0) 33.0(7.0) 9.0(1.0) 13.0(0.0) 0.0 (11.0) 23.0(3.5) 4.5 Refinement (16.0) 44.0(0.0) 12.0(2.0) 26.0(0.0) 0.0(9.0) 35.0(0.0) 6.0 Ours(9.0) 39.0(2.0) 19.0(1.0) 19.0(0.0) 0.0(5.0) 29.0(1.0) 9.5 GPT-5.2 Pass@k(24.0) 36.0(9.0) 19.0(2.0) 13.0(1.0) 3.0 (13.0) 24.5 (5.0) 11.0 Refinement (19.0) 88.0(7.0) 41.0(2.0) 55.0(1.0) 3.0 (10.5) 71.5 (4.0) 22.0 Codex(34.0) 82.0 (16.0) 70.0 (16.0) 84.0 (0.0) 22.0 (25.0) 83.0 (8.0) 46.0 Ours(20.0) 97.0(7.0) 90.0 (2.0) 100.0 (1.0) 76.0 (11.0) 98.5 (4.0) 83.0 We apply PopArt-style online normalizationˆr 2,t = (r 2,t − μ 2 )/σ 2 using running estimates(μ 2 ,σ 2 ). We update Q 2 for both the start pointp t and retrieved entitiesz ∈ c t using Eq. 4 withr = ˆr 2,t . When a refined kernel is feasible, as indicated byg feas (o t ) = 1, we store the kernel together with verifier feedback in memory for future retrieval and add it to the start setP(x) to expand the refinement search space. 3.5MULTI-GATE VERIFICATION The verifierVacts as the environment interface, providing robust feedback to guide the RL process. Given a task x and a generated kernel y t , it returns a structured outcome o t = V (x,y t ) = (g hack ,g comp ,g corr ,ℓ lat ),(7) whereg hack ,g comp ,g corr ∈0, 1denote the anti-hacking, compilation, and correctness gates, andℓ lat ∈ R + is the measured latency. A kernel is deemed feasible if and only if: g feas (o t )≜ g hack ∧ g comp ∧ g corr . Anti-hacking (g hack ). We implement a two-tier screening process. A rule-based filter first rejects trivial exploits (e.g., using high-leveltorchAPIs or constant-folding shortcuts). Survivors undergo a model-based inspection to identify subtle harness manipulations. Compilation (g comp ) & Correctness (g corr ). We verify successful compilation under the backend-specific toolchain, instantiated in our main study with the Ascend C toolchain. Correctness is validated by comparing outputs against the PyTorch reference:∥out y (x)− ref(x)∥ ≤ τ. The verifier provides fine-grained feedback, including mismatch localization and shape errors (details in Appendix F). Latency (ℓ lat ). For feasible kernels, we measure on-device execution time using backend-native profiling tools. In the primary Ascend setting, we use msprof and report the mean wall time across 3 profiling passes (Pipe, Memory, Resource) after warm-up. In extended experiments on CUDA, the same loop is instantiated with GPU-native profiling signals. 6 4EXPERIMENT 4.1EXPERIMENTAL SETUP Benchmark and Execution. We evaluate on L1 and L2 operators from KernelBench (Ouyang et al., 2025). Since KernelBench does not natively support Ascend C, we implement a compilation, deployment, and execution pipeline that maintains full compatibility with KernelBench PyTorch references while enabling the model to generate complete Ascend operator projects. Budget and metric. We enforce a strict per-operator budget ofT = 30iterations across all methods, encompassing both draft generation and iterative refinement. Functional correctness is verified with tolerances ofatol = rtol = 10 −2 . Our evaluation relies on three primary metrics: (i) Compilation Rate (CR), which measures the proportion of generated kernels that successfully compile, and (i) Correctness (Acc), which reports the percentage of operators for which a functionally valid solution is found within the budget.(i) Speedup measures the reduction in execution latency, defined asspeedup = L ref /L opt , whereL ref andL opt are the latencies of the reference and optimized kernels, respectively. Baselines. We compare EvoKernel against three baseline strategies using three models: Qwen3-Coder-30B-A3B- Instruct (Yang et al., 2025), DeepSeek-V3.2 (Liu et al., 2025), and GPT-5.2. Detailed configurations of these baselines can be found in Appendix D. • Pass@k: A stateless baseline generatingK = 30independent candidates per operator given a single demonstration. •Refinement: A stateful agentic loop that iteratively repairs compilation and correctness errors using verifier feedback. Upon finding a valid kernel, it transitions to hill-climbing for latency optimization, subject to a maximum budget of 30 iterations. •Codex by OpenAI: An autonomous agent based on GPT-5.2 with direct shell and file system access. It executes a “try-fail-evolve” loop, autonomously mutating the implementation based on execution logs until success or a budget of 30 verification attempts is exhausted. 4.2MAIN RESULTS We evaluate EvoKernel under a matched evaluation pipeline, focusing on compilation and correctness, as well as performance optimization after correctness. Compilation and correctness. Table 2 reports compilation rate (CR) and correctness (Acc) across two difficulty levels under a fixed budget T =30. EvoKernel achieves the strongest overall performance with GPT-5.2, reaching 98.5% CR and 83.0% Acc, substantially outperforming Codex (83.0% CR, 46.0% Acc) and Refinement (71.5% CR, 22.0% Acc). On Level 2, EvoKernel attains near-perfect compilation (100%) with 76% correctness. Despite Codex having autonomous shell and file system access, EvoKernel surpasses it by 15.5 points in CR and 37.0 points in Acc. On weaker backbones, the improvements are more moderate. EvoKernel achieves the highest Acc on both Qwen3-Coder-30B (5.5% vs. 4.0%) and DeepSeek-V3.2 (9.5% vs. 6.0%), with DeepSeek-V3.2 reaching 19% correctness on Level 1—more than doubling Pass@k. The Refinement baseline attains higher CR on DeepSeek- V3.2 (35.0% vs. 29.0%), suggesting that value-driven retrieval prioritizes generation quality over compilation attempts. Critically, Level 2 Acc remains at 0% for weaker models even when candidates compile (e.g., 19% CR on DeepSeek-V3.2), indicating that harder operators demand stronger generator capacity. Examining Round 1 through the final iteration reveals how effectively each method leverages the iterative process. On GPT-5.2, EvoKernel improves CR from 11.0% to 98.5% and Acc from 4.0% to 83.0%, representing an order- of-magnitude gain. In contrast, weaker models show limited improvement: Qwen3-Coder-30B increases Acc by +2.5 points, while DeepSeek-V3.2 improves by +8.5 points. This disparity reveals a key insight: the in-context learning capabilities of frontier LLMs prove critical for experience-driven approaches like ours. Crucially, it 7 does not weaken our method’s value; instead, it confirms that our agent is keeping pace with the cutting-edge advancements of base models. ActivationConvolutionFuseLossMatmulNormalizationPooling Category 0.0 0.2 0.4 0.6 0.8 1.0 Correctness Rate Correctness by Category (iter=30) 1.0× speedup 0.8-1.0× speedup 0.5-0.8× speedup 0.0-0.5× speedup Operators (n=159) 0 50 100 150 200 250 Speedup (first-correct / best) Speedup vs First Correct (iter=30) 1 2 3 4 102030 10 2 10 3 ms (1) Iter=26 60× 53_Min_reduction_over_a_dimension 51015202530 10 4 (2) Iter=29 13× 54_conv_standard_3D__square_input__square_kernel 102030 Iter 10 4 ms (3) Iter=30 13× 83_Conv3d_GroupNorm_Min_Clamp_Dropout 1015202530 Iter 200 400 600 800 (4) Iter=30 8× 82_conv_depthwise_2D_square_input_square_kernel Figure 2: Optimization outcomes. (Left) Category-level correctness and speedup distribution at budgetT =30; color segments show the fraction of correct kernels in each speedup tier relative to Torch-NPU. (Right) Within- operator speedup achieved by iterative refinement across 159 operators with≥1 valid optimization candidate beyond the initial correct draft; inset panels detail representative optimization trajectories. Optimization gains: within-operator speedup. Conditioned on reaching a correct draft, the refining stage further reduces latency. For each solved operator, we compare the initial draft, defined as the first feasible candidate, to the best candidate found within the remaining budget. This yields a median speedup of 3.60×, with an interquartile range of 1.38–10.05×. Although many operators remain slower than Torch-NPU (Figure 2), consistent within-operator gains indicate that the refinement process continues to improve performance beyond correctness. Figure 2 quantifies these gains across 159 operators with at least one valid optimization candidate beyond the initial correct draft. The distribution is long-tailed: while many operators exhibit modest improvements (s≈ 1–2×), a substantial subset benefits dramatically from continued optimization, with top performers achieving more than 200× speedup over their first correct version. Inset trajectories for four representative operators confirm that these gains emerge from systematic, incremental improvements across multiple iterations, rather than from single fortuitous generations. 4.3GENERALIZATION OF VALUE-DRIVEN MEMORY A core motivation for our memory design is reusability: high-utility past experiences should accelerate learning on subsequent ones. We verify this hypothesis by evaluating transfer across difficulty levels and generator backbones. 2468101214161820 0.0 0.2 0.4 0.6 0.8 1.0 Cumulative Correctness Rate 34.0% 53.0% 64.0% Iteration Level Mixing Strategy (L2 Operators) L2 ScratchL1+L2 MixedL1L2 DeepSeek-V3.2Qwen3-Coder 0 20 40 60 80 100 Correctness (%) 6% 4% 58% 32% (a) Correctness Rate DeepSeek-V3.2Qwen3-Coder 0 20 40 60 80 100 Compilation (%) 26% 14% 80% 84% (b) Compilation Rate Memory Transfer (50 Operators) w/o mem.w/ mem. Figure 3: Transfer and generalization. (Left) Transfer across difficulty levels: cumulative success rate on L2 under different stream compositions. (Right) Transfer across generator backbones: performance on held-out operators when reusing memory built with GPT-5.2. 8 Transfer across difficulty levels. We study whether memory accumulated on easier L1 operators transfers to harder L2 operators. We consider three setups: • L2 Scratch: agent iterates from scratch on L2 operators. • L1+L2 Mixed: the agent iterates from scratch on a mixed operator set containing both L1 and L2. •L1→L2: the agent first iterates on L1, then continues iterating on the L2 operator set initialized with the resulting L1 memory. In Figure 3 and Table 3, the L1→L2 stream exhibits the fastest warm-up and highest final performance. By iterationt = 17, it achieves 64% L2 correctness, outperforming L1+L2 Mixed (53%) by 11% and L2 Scratch (34%) by 30%. Crucially, the transfer allows the agent to solve its first L2 operator four iterations earlier than the scratch baseline. This confirms that foundational patterns learned from simpler tasks effectively bootstrap progress on harder problems. Table 3: Cross-level transfer summary on L2 at final iteration. SetupCR (%)Acc (%) L2 Scratch88.034.0 L1+L2 Mixed98.053.0 L1→L297.064.0 Transfer across generator backbones. We further assess whether memory constructed by a strong model (GPT-5.2) can improve the performance of weaker backbones (DeepSeek-V3.2, Qwen3-Coder-30B). We evaluate on a held-out set of 50 operators (30 L1, 20 L2), initializing the agent with a filtered GPT-5.2 memory bank where traces from the test operators are excluded to prevent leakage. Figure 3 (right) shows that the learned memory transfers well across generator backbones. For DeepSeek, adding memory improves compilation from 26% to 80% and correctness from 6% to 58%. For Qwen, memory yields a similarly large compilation gain (14%→84%) with a smaller but substantial correctness gain (4%→32%). Overall, memory appears to provide backbone-agnostic operator constraints and debugging cues that greatly reduce non-compiling attempts, while the remaining compilation–correctness gap (especially for Qwen) suggests semantic validity remains the dominant bottleneck. 4.4BEYOND KERNELBENCH AND CANN The main benchmark in this paper is Ascend C KernelBench, but an important question is whether the learned memory and refinement policy continue to help on workloads that fall outside this training distribution. To test this, we evaluate EvoKernel on the Attention Set operator suites and on mHC kernels (Xie et al., 2025) derived from recent DeepSeek architectures. Table 4: Initial scaling-out results beyond the main KernelBench study. WorkloadPlatform# OpsCR (%)Acc (%)Fast 1 Ratio (%) Attention SetCUDA70100.097.172.1 Attention SetAscend70100.078.621.8 KernelBenchCUDA250100.0100.068.0 mHC Kernels (DeepSeek)Ascend1586.766.760.0 Attention Set operators. On CUDA, EvoKernel scales cleanly from the KernelBench-style setting to the Attention Set workloads. On the 70-operator Attention Set (excluding non-attention operators such as GEMM, RoPE, and Router), the system reaches 100% compilation and 97.1% correctness after 30 outer iterations. On the same 9 Attention Set on CUDA with KernelBench operators included, EvoKernel achieves 100% compilation and 100% correctness on all 250 operators. These results indicate that the memory mechanism transfers beyond the operator families emphasized in the main benchmark and remains effective on more application-driven kernels. Figure 7 shows the optimization timeline and performance comparison for the CUDA Attention Set. Ascend Attention Set and DeepSeek mHC kernels. More importantly for the cold-start setting emphasized in this paper, the same methodology also transfers to new Ascend C workloads outside the original KernelBench distribu- tion. On the 70-operator Ascend Attention Set, EvoKernel reaches 100.0% compilation and 78.6% correctness after 30 iterations. We further evaluate on 15 mHC kernels targeting a recent DeepSeek architectural motif on Ascend (CANN 8.5.0). EvoKernel obtains 10 correct implementations, and 6 of these outperform the PyTorch baseline. Representative wins includeSinkhornKnoppwith 41.96×speedup,OrthostochasticProjectwith 2.94×, andMhcPostBlockwith 2.88×. Figure 4 shows the optimization timeline and per-operator performance for all 15 mHC kernels over 30 iterations (merged across three experiment series). 123456789101112131415161718192021222324252627282930 Outer Iteration 01_SinkhornKnopp 05_StreamWrite 04_StreamMix 06_MhcUpdate 03_StreamWeightedSum 15_MhcPostBlock 10_OrthostochasticProject 02_MhcProjector 14_MhcPreBlock 12_OptimizedMHCLayerWithFusion 0.6740.060.040.0340.027 0.1850.0590.0420.037 0.6880.2670.0620.028 35.190.7140.3090.250.0720.0490.04 0.2080.1520.0360.0180.015 233629.5124.510.270.714 0.0770.0480.0410.036 217.575.0856.715.029.1377.099 3567112.6 22.34 mHC Kernels (Ascend) Correctness & Performance Optimization Timeline First solveBest perf improvedBest performance 642024 Relative Performance log(Baseline / Agent) 01_SinkhornKnopp 02_MhcProjector 03_StreamWeightedSum 04_StreamMix 05_StreamWrite 06_MhcUpdate 10_OrthostochasticProject 12_OptimizedMHCLayerWithFusion 14_MhcPreBlock 15_MhcPostBlock 41.96x 0.36x 2.67x 1.48x 0.35x 1.47x 2.94x 0.06x 0.01x 2.88x Best Correct Run vs PyTorch Baseline (mHC Kernels (Ascend)) PyTorch Baseline (1.0x) 0.00x0.02x0.06x0.25x1.00x4.00x16.00x64.00x Equivalent Speedup (Baseline / Agent) Figure 4: mHC Kernels (Ascend): Optimization timeline and performance vs. Torch-NPU baseline for 15 DeepSeek mHC operators over 30 iterations (merged across three experiment series). (Left) Correctness and performance optimization timeline. (Right) Best correct run vs. baseline in log 2 speedup. Taken together, these scaling-out results suggest that EvoKernel is not simply memorizing KernelBench operator templates. Instead, the framework appears able to reuse memory and profiling-guided refinement to adapt to both new operator families and new architectural motifs, while still preserving the paper’s primary emphasis on data-scarce Ascend C kernel generation. 4.5ABLATIONS 4.5.1VALUE-DRIVEN VERSUS HEURISTIC-DRIVEN RETRIEVAL We assess the impact of learned value estimates by comparing our full value-driven pipeline against a heuristic- driven variant. Both settings use the L1→L2 transfer protocol (Section 4.3) and run for 30 L2 iterations per operator, inheriting the same L1 memory. The only difference lies in the selection mechanism: • Value-Driven (Ours): Selects context and optimization start points using ε-greedy over learned Q-values. •Heuristic-Driven: Selects context based solely on semantic similarity and chooses optimization start points based on the highest historical performance. Figure 5 (left) tracks cumulative correctness and compilation rates. While both methods perform similarly in the early stages (reaching 48% correctness by iteration 14), the value-driven approach diverges significantly thereafter. 10 By iteration 30, it achieves 77% correctness and 100% compilation, compared to 67% and 97% for the heuristic baseline. This indicates that while heuristics suffice for initial bootstrapping, learned value estimates provide a crucial exploitation signal for solving the long tail of difficult operators. 51015202530 Iteration 0.0 0.2 0.4 0.6 0.8 1.0 Cumulative Correctness Rate 67.0% 77.0% Iter=14 (a) Correctness Rate 51015202530 Iteration 0.0 0.2 0.4 0.6 0.8 1.0 Cumulative Compilation Rate 97.0% 100.0% Iter=5 (b) Compilation Rate Q Value Ablation (L2 Operators) w/o Q valuew/ Q value 510152025 Iteration 0.0 0.2 0.4 0.6 0.8 1.0 Cumulative Correctness Rate 84.0% 67.0% Iter=24 Top-K Similarity Search Observation (L1 Operators) CorrectnessCompilation 0.0 0.2 0.4 0.6 0.8 1.0 Cumulative Compilation Rate 98.0% Figure 5: Retrieval ablations. (Left) Value-driven vs. heuristic retrieval on L2 operators (same L1 memory and ε-greedy schedule). (Right) Effect of increasing retrieval pool sizeKat iteration 24; cumulative correctness and compilation rates on L1 operators. 4.5.2MULTI-TASK MEMORY SHARING VERSUS PER-TASK REFINEMENT To isolate the contribution of cross-task memory sharing, we compare EvoKernel against the Refinement baseline under identical per-operator iteration budgets (Table 2). Refinement can be viewed as a degenerate instance of our framework: restricting the memory bank to a single operator eliminates cross-task retrieval, reducing the agent to iterative self-refinement. This controlled ablation thus directly quantifies the benefit of a global, shared memory bank over per-task isolated iteration. Results reveal that cross-task sharing yields substantial gains, particularly on Level 2 operators. With GPT-5.2, EvoKernel raises the Level 2 compilation rate from 55.0% to 100.0% and accuracy from 3.0% to 76.0%. Level 1 exhibits more moderate improvements (+9 p CR, +49 p Acc). These findings indicate that, although within- operator refinement provides a useful signal, the ability to transfer experience across tasks confers additional, complementary benefits that isolated iteration cannot achieve. 4.6DISCUSSION Explicit versus Emergent Curricula. Our results demonstrate that value-driven memory induces adaptive curriculum learning without explicit task ordering. When we impose an explicit L1→L2 curriculum (Table 3), the agent benefits from a warm start, as L1 memory acts as foundational scaffolding that accelerates early L2 progress despite the complexity gap. Crucially, however, even under a L1+L2 Mixed setting with no prescribed ordering, the retrieval policy autonomously reconstructs a soft curriculum. Figure 6 exemplifies this emergent behavior for36RMSNormwithin the mixed setting: the agent first solves simpler operators, which then serve as retrieved references to facilitate the solution of harder ones, naturally forming a dependency chain without manual intervention. Scaling to out-of-distribution workloads. The additional results in Table 4 strengthen this interpretation. The framework transfers not only across difficulty levels within KernelBench, but also to workloads that differ materially from the main training distribution: the Attention Set and recent DeepSeek mHC kernels. In particular, the Ascend mHC results indicate that the system can reuse accumulated memories to tackle new architectural motifs rather than only variants of benchmark operators. The Ascend Attention Set results point in the same direction, with 78.6% correctness on 70 operators, suggesting that the agent’s benefit persists even when the workload shifts toward application-driven kernels. Why value-driven memory outperforms stateless baselines. Pass@ksampling treats each generation indepen- dently, forfeiting any cross-attempt learning. Iterative refinement (e.g., Codex) accumulates feedback within a single operator but discards it afterward, preventing cross-operator transfer. In contrast, our approach persists and 11 S1=iter1S2=iter2S3=iter4S4=iter6S5=iter7S6=iter12 22_Tanh ACTIVATION 29_Softplus ACTIVATION 23_Softmax ACTIVATION 34_InstanceNorm NORMALIZATION 37_FrobeniusNorm_ NORMALIZATION 36_RMSNorm_ NORMALIZATION Experience Transfer for 36_RMSNorm_ SolvedUnsolvedExample used Figure 6: Experience transfer dependency graph of36RMSNorm. Arrows trace causal references at first-solve iterations, revealing an emergent curriculum from simple to complex operators. values experiences across both attempts and tasks, enabling the agent to bootstrap harder problems from easier ones and to amortize debugging effort across the entire operator population. Impact of candidate pool size. In our experiments, the candidate pool size|C(x)|is controlled by a multiplierλ applied to the final retrieval countN. Smaller candidate pools risk missing valuable context, while larger ones may introduce noise and dilute the signal from high-value entries. Initially, we setλ = 2, resulting in a convergence point with 67% correctness. Upon increasingλby a factor of 15, correctness improved sharply to 84% by iteration 26. This suggests that dynamically expanding the candidate pool during training allows the Q-value policy to discover previously overlooked high-utility entries. The optimal multiplier remains an important area for future research, as there is likely a sweet spot that balances coverage and efficiency. In our experiments, we observed that gradually increasingλallowed for a controlled replacement of context, ultimately improving model performance. 5CONCLUSION AND FUTURE WORK We presented EvoKernel, a value-driven memory agent addressing cold-start kernel synthesis by learning stage- specific Q-values for retrieval over a self-evolving memory bank. A central insight is that frontier LLMs have enhanced in-context learning capabilities, enabling effective generalization from retrieved demonstrations even in cold-start kernel synthesis scenarios. This emergent ability makes memory-based, non-parametric approaches practically viable. Our additional scaling results on the Attention Set and recent DeepSeek mHC kernels further suggest that the learned memory is not confined to the original KernelBench distribution. More broadly, the value-driven memory paradigm may benefit other cold-start domains with binary verification signals, and we anticipate that as LLMs continue to improve, memory-augmented approaches will enable autonomous mastery of an ever-wider range of specialized tasks. Beyond technical gains, these results suggest value-driven memory can democratize data-scarce programming expertise (e.g., NPU kernel synthesis), helping bridge expert shortages as hardware diversifies and pointing toward AI systems that adapt to new domains with minimal data. Potential future work includes extending the framework to other emerging DSLs to verify cross-architecture universality, exploring knowledge distillation to reduce reliance on large commercial models, and incorporating denser reward signals to improve sample efficiency. 12 REFERENCES Scott Barnett, Stefanus Kurniawan, Srikanth Thudumu, Zach Brannelly, and Mohamed Abdelrazek. Seven failure points when engineering a retrieval augmented generation system. In Proceedings of the IEEE/ACM 3rd International Conference on AI Engineering-Software Engineering for AI, p. 194–199, 2024. Carlo Baronio, Pietro Marsella, Ben Pan, Simon Guo, and Silas Alberti. Kevin: Multi-turn RL for generating CUDA kernels. arXiv preprint arXiv:2507.11948, 2025. Meng Cao, Lei Shu, Lei Yu, Yun Zhu, Nevan Wichers, Yinxiao Liu, and Lei Meng. Beyond sparse rewards: Enhanc- ing reinforcement learning with language model critique in text generation. arXiv preprint arXiv:2401.07382, 2024. P Chhikara, D Khant, S Aryan, T Singh, and D Yadav. Mem0: Building production-ready AI agents with scalable long-term memory (2025). URL https://arxiv.org/abs/2504.19413, 2025. Jack Choquette, Wishwesh Gandhi, Olivier Giroux, Nick Stam, and Ronny Krashinsky. NVIDIA A100 tensor core GPU: Performance and innovation. IEEE Micro, 41(2):29–35, 2021. Hyung Won Chung, Le Hou, Shayne Longpre, Barret Zoph, Yi Tay, William Fedus, Yunxuan Li, Xuezhi Wang, Mostafa Dehghani, Siddhartha Brahma, et al. Scaling instruction-finetuned language models. Journal of Machine Learning Research, 25(70):1–53, 2024. Emile Contal and Garrin McGoldrick. RAGSys: Item-cold-start recommender as RAG system. arXiv preprint arXiv:2405.17587, 2024. Juncheng Dong, Yang Yang, Tao Liu, Yang Wang, Feng Qi, Vahid Tarokh, Kaushik Rangadurai, and Shuang Yang. STARK: Strategic team of agents for refining kernels. arXiv preprint arXiv:2510.16996, 2025. Jinye Du, Quan Yuan, Zuyao Zhang, Yanzhi Yi, Jiahui Hu, Wangyi Chen, Yiyang Zhu, Qishui Zheng, Wenxiang Zou, Xiangyu Chang, et al. AKG kernel agent: A multi-agent framework for cross-platform kernel synthesis. arXiv preprint arXiv:2512.23424, 2025. Runnan Fang, Yuan Liang, Xiaobin Wang, Jialong Wu, Shuofei Qiao, Pengjun Xie, Fei Huang, Huajun Chen, and Ningyu Zhang. MEMP: Exploring agent procedural memory. arXiv preprint arXiv:2508.06433, 2025. Matteo Hessel, Hubert Soyer, Lasse Espeholt, Wojciech Czarnecki, Simon Schmitt, and Hado Van Hasselt. Multi-task deep reinforcement learning with PopArt. In Proceedings of the AAAI Conference on Artificial Intelligence, volume 33, p. 3796–3803, 2019. Gautier Izacard, Patrick Lewis, Maria Lomeli, Lucas Hosseini, Fabio Petroni, Timo Schick, Jane Dwivedi-Yu, Armand Joulin, Sebastian Riedel, and Edouard Grave. Atlas: Few-shot learning with retrieval augmented language models. Journal of Machine Learning Research, 24(251):1–43, 2023. Naman Jain, King Han, Alex Gu, Wen-Ding Li, Fanjia Yan, Tianjun Zhang, Sida Wang, Armando Solar-Lezama, Koushik Sen, and Ion Stoica. LiveCodeBench: Holistic and contamination free evaluation of large language models for code. arXiv preprint arXiv:2403.07974, 2024. Sathvik Joel, Jie JW Wu, and Fatemeh H Fard. A survey on LLM-based code generation for low-resource and domain-specific programming languages, 2024. URL https://arxiv. org/abs/2410.03981, 2024. Norm Jouppi, George Kurian, Sheng Li, Peter Ma, Rahul Nagarajan, Lifeng Nai, Nishant Patil, Suvinay Subra- manian, Andy Swing, Brian Towles, et al. TPU v4: An optically reconfigurable supercomputer for machine learning with hardware support for embeddings. In Proceedings of the 50th annual international symposium on computer architecture, p. 1–14, 2023. Sham Machandranath Kakade. On the sample complexity of reinforcement learning. University of London, University College London (United Kingdom), 2003. 13 Jared Kaplan, Sam McCandlish, Tom Henighan, Tom B Brown, Benjamin Chess, Rewon Child, Scott Gray, Alec Radford, Jeffrey Wu, and Dario Amodei. Scaling laws for neural language models. arXiv preprint arXiv:2001.08361, 2020. Aida Kostikova, Zhipin Wang, Deidamea Bajri, Ole P ̈ utz, Benjamin Paaßen, and Steffen Eger. LLLMs: A data- driven survey of evolving research on limitations of large language models. arXiv preprint arXiv:2505.19240, 2025. Patrick Lewis, Ethan Perez, Aleksandra Piktus, Fabio Petroni, Vladimir Karpukhin, Naman Goyal, Heinrich K ̈ uttler, Mike Lewis, Wen-tau Yih, Tim Rockt ̈ aschel, et al. Retrieval-augmented generation for knowledge- intensive NLP tasks. Advances in Neural Information Processing Systems, 33:9459–9474, 2020. Jianling Li, Shangzhan Li, Zhenye Gao, Qi Shi, Yuxuan Li, Zefan Wang, Jiacheng Huang, WangHaojie WangHao- jie, Jianrong Wang, Xu Han, et al. TritonBench: Benchmarking large language model capabilities for generating Triton operators. In Findings of the Association for Computational Linguistics: ACL 2025, p. 23053–23066, 2025a. Shangzhan Li, Zefan Wang, Ye He, Yuxuan Li, Qi Shi, Jianling Li, Yonggang Hu, Wanxiang Che, Xu Han, Zhiyuan Liu, et al. AutoTriton: Automatic Triton programming with reinforcement learning in LLMs. arXiv preprint arXiv:2507.05687, 2025b. Zhiyu Li, Chenyang Xi, Chunyu Li, Ding Chen, Boyu Chen, Shichao Song, Simin Niu, Hanyu Wang, Jiawei Yang, Chen Tang, et al. Memos: A memory OS for AI system. arXiv preprint arXiv:2507.03724, 2025c. Heng Liao, Jiajin Tu, Jing Xia, Hu Liu, Xiping Zhou, Honghui Yuan, and Yuxing Hu. Ascend: A scalable and unified architecture for ubiquitous deep neural network computing: Industry track paper. In 2021 IEEE International Symposium on High-Performance Computer Architecture (HPCA), p. 789–801. IEEE, 2021. Aixin Liu, Aoxue Mei, Bangcai Lin, Bing Xue, Bingxuan Wang, Bingzheng Xu, Bochao Wu, Bowei Zhang, Chaofan Lin, Chen Dong, et al. DeepSeek-v3.2: Pushing the frontier of open large language models. arXiv preprint arXiv:2512.02556, 2025. Shayne Longpre, Le Hou, Tu Vu, Albert Webson, Hyung Won Chung, Yi Tay, Denny Zhou, Quoc V Le, Barret Zoph, Jason Wei, et al. The FLAN collection: Designing data and methods for effective instruction tuning. In International Conference on Machine Learning, p. 22631–22648. PMLR, 2023. Aman Madaan, Niket Tandon, Peter Clark, and Yiming Yang. MemPrompt: Memory-assisted prompt editing with user feedback, 2022. Aman Madaan, Niket Tandon, Prakhar Gupta, Skyler Hallinan, Luyu Gao, Sarah Wiegreffe, Uri Alon, Nouha Dziri, Shrimai Prabhumoye, Yiming Yang, et al. Self-Refine: Iterative refinement with self-feedback. Advances in Neural Information Processing Systems, 36:46534–46594, 2023. Shervin Minaee, Tomas Mikolov, Narjes Nikzad, Meysam Chenaghlu, Richard Socher, Xavier Amatriain, and Jianfeng Gao. Large language models: A survey. arXiv preprint arXiv:2402.06196, 2024. Alexander Novikov, Ng ˆ an V ̃ u, Marvin Eisenberger, Emilien Dupont, Po-Sen Huang, Adam Zsolt Wagner, Sergey Shirobokov, Borislav Kozlovskii, Francisco JR Ruiz, Abbas Mehrabian, et al. AlphaEvolve: A coding agent for scientific and algorithmic discovery. arXiv preprint arXiv:2506.13131, 2025. Anne Ouyang, Simon Guo, Simran Arora, Alex L Zhang, William Hu, Christopher Re, and Azalia Mirhoseini. KernelBench: Can LLMs write efficient GPU kernels? In Aarti Singh, Maryam Fazel, Daniel Hsu, Simon Lacoste-Julien, Felix Berkenkamp, Tegan Maharaj, Kiri Wagstaff, and Jerry Zhu (eds.), Proceedings of the 42nd International Conference on Machine Learning, volume 267 of Proceedings of Machine Learning Research, p. 47356–47415. PMLR, 13–19 Jul 2025. Charles Packer, Vivian Fang, Shishir G. Patil, Kevin Lin, Sarah Wooders, and Joseph Gonzalez. MemGPT: Towards LLMs as operating systems. ArXiv, abs/2310.08560, 2023. 14 Joon Sung Park, Joseph O’Brien, Carrie Jun Cai, Meredith Ringel Morris, Percy Liang, and Michael S Bernstein. Generative agents: Interactive simulacra of human behavior. In Proceedings of the 36th annual acm symposium on user interface software and technology, p. 1–22, 2023. Han Qi, Haochen Yang, Qiaosheng Zhang, and Zhuoran Yang. Sample-efficient reinforcement learning from human feedback via information-directed sampling. arXiv preprint arXiv:2502.05434, 2025. Dezhi Ran, Shuxiao Xie, Mingfang Ji, Ziyue Hua, Mengzhou Wu, Yuan Cao, Yuzhe Guo, Yu Hao, Linyi Li, Yitao Hu, et al. KernelBand: Boosting LLM-based kernel optimization with a hierarchical and hardware-aware multi-armed bandit. arXiv preprint arXiv:2511.18868, 2025. Herbert Robbins and David Siegmund. A convergence theorem for non negative almost supermartingales and some applications. In Optimizing methods in statistics, p. 233–257. Elsevier, 1971. Noah Shinn, Federico Cassano, Ashwin Gopinath, Karthik Narasimhan, and Shunyu Yao. Reflexion: Language agents with verbal reinforcement learning. Advances in Neural Information Processing Systems, 36:8634–8652, 2023. Cristina Silvano, Daniele Ielmini, Fabrizio Ferrandi, Leandro Fiorin, Serena Curzel, Luca Benini, Francesco Conti, Angelo Garofalo, Cristian Zambelli, Enrico Calore, et al. A survey on deep learning hardware accelerators for heterogeneous HPC platforms. ACM Computing Surveys, 57(11):1–39, 2025. Richard S Sutton and Andrew G Barto. Reinforcement learning: An introduction. MIT press, 2018. Guanzhi Wang, Yuqi Xie, Yunfan Jiang, Ajay Mandlekar, Chaowei Xiao, Yuke Zhu, Linxi Fan, and Anima Anand- kumar. Voyager: An open-ended embodied agent with large language models. arXiv preprint arXiv:2305.16291, 2023. Xinyi Wang, Antonis Antoniades, Yanai Elazar, Alfonso Amayuelas, Alon Albalak, Kexun Zhang, and William Wang. Generalization vs. memorization: Tracing language models’ capabilities back to pretraining data. In Y. Yue, A. Garg, N. Peng, F. Sha, and R. Yu (eds.), International Conference on Learning Representations, volume 2025, p. 49948–49968, 2025. Zhongzhen Wen, Yinghui Zhang, Zhong Li, Zhongxin Liu, Linna Xie, and Tian Zhang. MultiKernelBench: A multi-platform benchmark for kernel generation. arXiv eprints, p. arXiv–2507, 2025. Jiin Woo, Shaowei Zhu, Allen Nie, Zhen Jia, Yida Wang, and Youngsuk Park. TritonRL: Training LLMs to think and code Triton without cheating. arXiv preprint arXiv:2510.17891, 2025. Peng Wu. PyTorch 2.0: The journey to bringing compiler technologies to the core of PyTorch (keynote). In Proceedings of the 21st ACM/IEEE International Symposium on Code Generation and Optimization, p. 1–1, 2023. Rong Wu, Xiaoman Wang, Jianbiao Mei, Pinlong Cai, Daocheng Fu, Cheng Yang, Licheng Wen, Xuemeng Yang, Yufan Shen, Yuxin Wang, et al. Evolver: Self-evolving LLM agents through an experience-driven lifecycle. arXiv preprint arXiv:2510.16079, 2025. Zhenda Xie, Yixuan Wei, Huanqi Cao, Chenggang Zhao, Chengqi Deng, Jiashi Li, Damai Dai, Huazuo Gao, Jiang Chang, Kuai Yu, et al. mhc: Manifold-constrained hyper-connections. arXiv preprint arXiv:2512.24880, 2025. Weixiang Yan, Haitian Liu, Yunkun Wang, Yunzhe Li, Qian Chen, Wen Wang, Tingyu Lin, Weishan Zhao, Li Zhu, Hari Sundaram, et al. CodeScope: An execution-based multilingual multitask multidimensional benchmark for evaluating LLMs on code understanding and generation. In Proceedings of the 62nd Annual Meeting of the Association for Computational Linguistics (Volume 1: Long Papers), p. 5511–5558, 2024. An Yang, Anfeng Li, Baosong Yang, Beichen Zhang, Binyuan Hui, Bo Zheng, Bowen Yu, Chang Gao, Chengen Huang, Chenxu Lv, et al. Qwen3 technical report. arXiv preprint arXiv:2505.09388, 2025. 15 Shunyu Yao, Dian Yu, Jeffrey Zhao, Izhak Shafran, Tom Griffiths, Yuan Cao, and Karthik Narasimhan. Tree of thoughts: Deliberate problem solving with large language models. Advances in Neural Information Processing Systems, 36:11809–11822, 2023. Yang Yu, Peiyu Zang, Chi Hsu Tsai, Haiming Wu, Yixin Shen, Jialing Zhang, Haoyu Wang, Zhiyou Xiao, Jingze Shi, Yuyu Luo, et al. Towards automated kernel generation in the era of LLMs. arXiv preprint arXiv:2601.15727, 2026. Shengtao Zhang, Jiaqian Wang, Ruiwen Zhou, Junwei Liao, Yuchen Feng, Weinan Zhang, Ying Wen, Zhiyu Li, Feiyu Xiong, Yutao Qi, et al. MemRL: Self-evolving agents via runtime reinforcement learning on episodic memory. arXiv preprint arXiv:2601.03192, 2026. Zihan Zhang, Yuxin Chen, Jason Lee, and Simon S Du. Settling the sample complexity of online reinforcement learning. Journal of the ACM, 72(3):1–63, 2025. Chunting Zhou, Pengfei Liu, Puxin Xu, Srinivasan Iyer, Jiao Sun, Yuning Mao, Xuezhe Ma, Avia Efrat, Ping Yu, Lili Yu, et al. LIMA: Less is more for alignment. Advances in Neural Information Processing Systems, 36: 55006–55021, 2023. Huichi Zhou, Yihang Chen, Siyuan Guo, Xue Yan, Kin Hei Lee, Zihan Wang, Ka Yiu Lee, Guchun Zhang, Kun Shao, Linyi Yang, and Jun Wang. Memento: Fine-tuning LLM agents without fine-tuning LLMs. arXiv preprint arXiv: 2508.16153, 2025. Xinguo Zhu, Shaohui Peng, Jiaming Guo, Yunji Chen, Qi Guo, Yuanbo Wen, Hang Qin, Ruizhi Chen, Qirui Zhou, Ke Gao, et al. QiMeng-Kernel: Macro-thinking micro-coding paradigm for LLM-based high-performance GPU kernel generation. arXiv preprint arXiv:2511.20100, 2025. 16 APROOFS FOR VALUE UPDATE STABILITY AND CONVERGENCE This appendix establishes theoretical guarantees for the value-driven memory system introduced in Section 3.2. We prove three results: (1) boundedness of value iterates under bounded rewards, (2) stability of online reward normalization, and (3) convergence of the bandit-style update rule. Together, these lemmas ensure that the retrieval policy remains well-behaved throughout the agent’s lifetime. A.1NOTATION AND SETUP Fix a memory entryiand a stages∈draft,optimize. Each time entryiis retrieved, the system observes a scalar reward R t . The bandit-style value update is Q t+1 = Q t + α t (R t − Q t ) = (1− α t )Q t + α t R t , α t ∈ (0, 1].(8) This is the standard incremental mean estimator used throughout reinforcement learning (Sutton & Barto, 2018). Reward definitions by stage. (i) Draft stage: Binary rewardR t ∈+1,−1based on feasibility. (i) Optimize stage: Given speedup ratioρ t > 0, the raw reward isr raw,t = tanh(logρ t ) ∈ (−1, 1). Optionally, we apply z-score normalization: R t = (r raw,t − μ t−1 )/σ t−1 . A.2BOUNDEDNESS OF VALUE ITERATES Lemma 1 (Bounded Rewards Imply Bounded Values). Suppose|R t | ≤ R max for alltalmost surely, and α t ∈ (0, 1]. If Q 0 ∈ [−R max ,R max ], then Q t ∈ [−R max ,R max ] for all t. Proof. By induction. The updateQ t+1 = (1− α t )Q t + α t R t is a convex combination ofQ t andR t . If both lie in [−R max ,R max ], so does Q t+1 . Corollary 2 (Boundedness of Raw Optimization Reward). For anyρ t > 0, we haver raw,t = tanh(logρ t ) ∈ (−1, 1). Proof. Since ρ t > 0, logρ t ∈ R, and tanh : R→ (−1, 1). Remark 3 (Z-Score Normalization Requires Safeguards). The z-score transformationR t = (r raw,t − μ t−1 )/σ t−1 can be unbounded whenσ t−1 → 0. We ensure boundedness via either: (i) a variance floorˆσ t−1 := maxσ t−1 ,σ min , yielding|R t |≤ 2/σ min ; or (i) output clipping R t := clip(R t ;−B,B). Remark 4 (Error Clipping Alone Is Insufficient). An alternative updateQ t+1 = Q t + α· clip(R t − Q t ;−C,C) bounds the per-step change but not the iterates themselves. IfR t ≡ M ≫ Q 0 , thenQ t = Q 0 + tαC → ∞. Hence, reward boundedness (Lemma 1) is essential. A.3STABILITY OF ONLINE NORMALIZATION Lemma 5 (Convergence of Running Statistics). Letr raw,t t≥1 be a strictly stationary ergodic process with E[r 2 raw,1 ] <∞ and Var(r raw,1 ) = σ 2 > 0. Define μ t := 1 t t X k=1 r raw,k , σ t := v u u t 1 t t X k=1 (r raw,k − μ t ) 2 .(9) Thenμ t → μ := E[r raw,1 ]andσ t → σalmost surely. Moreover, the normalization mapf t (r) := (r− μ t )/σ t converges uniformly on bounded sets to f ∞ (r) := (r− μ)/σ. Proof.By the ergodic theorem,μ t → μa.s. Writingσ 2 t = 1 t P k r 2 raw,k − μ 2 t and applying ergodicity to both terms givesσ 2 t → σ 2 a.s. Continuity of √ ·on(0,∞)yieldsσ t → σ. For uniform convergence on a bounded set J : |f t (r)− f ∞ (r)|≤ |μ− μ t | σ t +|r− μ|· 1 σ t − 1 σ → 0 uniformly on J since σ t → σ > 0. 17 Remark 6 (Relation to PopArt). PopArt (Hessel et al., 2019) rescales network outputs when(μ,σ)change to preserve unnormalized predictions. Our scheme omits this rescaling; Lemma 5 shows the weaker but sufficient result that the normalization map stabilizes asymptotically. A.4CONVERGENCE OF THE BANDIT UPDATE We analyze two regimes: constant step size (tracking) and decreasing step size (convergence). Lemma 7 (Constant Step Size: EMA Dynamics). LetR t be i.i.d. with meanμand varianceσ 2 R <∞. Under constant α∈ (0, 1): (i) E[Q t ]→ μ as t→∞. (i) Var(Q t )→ α 2−α σ 2 R = O(α) for small α. (i) Q t converges in distribution to a unique stationary distribution centered at μ. Proof. Unrolling the recursion: Q t = (1− α) t Q 0 + α P t−1 k=0 (1− α) t−1−k R k . • Mean: E[Q t ] = (1− α) t Q 0 + μ(1− (1− α) t )→ μ. • Variance: Var(Q t ) = (1− α) 2t Var(Q 0 ) + α 2 σ 2 R P t−1 j=0 (1− α) 2j → α 2−α σ 2 R . •Distribution: The recursion defines an affine iterated function system with contraction(1− α) < 1, implying geometric ergodicity (Sutton & Barto, 2018). Lemma 8 (Decreasing Step Size: Almost Sure Convergence). Assume|R t |≤ R max a.s.,E[R t |F t ] = μ, andα t satisfies the Robbins-Monro conditions: P t α t =∞ and P t α 2 t <∞. Then Q t → μ almost surely. Proof.Definee t := Q t − μandξ t+1 := R t − μ. Thene t+1 = (1− α t )e t + α t ξ t+1 . LetV t := e 2 t . By direct computation: E[V t+1 |F t ]≤ (1− α t ) 2 V t + α 2 t σ 2 ξ ≤ V t − α t V t + α 2 t σ 2 ξ . By the Robbins-Siegmund theorem (Robbins & Siegmund, 1971),V t converges a.s. and P t α t V t <∞ . Since P t α t =∞, we must have V t → 0 a.s., hence Q t → μ. A.5SUMMARY The three results work in concert: Lemma 1 ensures value iterates remain in a safe range when rewards are bounded (which Corollary 2 and Remark 3 guarantee for our reward definitions). Lemma 5 ensures the normalization map stabilizes over time. Finally, Lemmas 7 and 8 establish that the value estimates track (constantα) or converge to (decreasingα t ) the true expected utility. Together, these guarantees ensure stable, well-behaved retrieval throughout the agent’s lifetime. Table 5: Operators where EvoKernel outperforms Torch-NPU in latency. OperatorTorch-NPU msEvoKernel msSpeedup 32HardTanh23.8734.1995.69× 30Softsign34.7569.5983.62× 45AveragePooling2D3814.7231725.4432.21× 20 LeakyReLU9.5259.5111.00× 18 BOPERATOR-LEVEL PERFORMANCE RESULTS This appendix provides select operator examples demonstrating performance comparisons. To contextualize absolute performance, we normalize latency by Torch-NPU and compare against other Ascend C approaches under the same MultiKernelBench harness. Table 5 lists example operators where EvoKernel outperforms Torch-NPU. CVERIFICATION STAGE “ANTI-HACKING” SCREENING In the context of this work, “anti-hacking” refers to the architectural enforcement of the Ascend C programming paradigm. It is designed to prevent a generated solution from bypassing the intended custom operator path by re-implementing semantics in Python (withinmodelsrc) or in the PyTorch binding glue (pythonbindsrc), rather than putting the computational logic into the Ascend C kernel (kernelsrc) and host tiling code. The verification subsystem implements this as a two-layer audit: 1.Rule-based screening (Static/Deterministic): Hard rules that reject common “semantic bypass” pat- terns. 2. Model-based screening (LLM Auditor): A prompt-driven judgment of “architectural integrity” that detects subtle bypass patterns not covered by static rules. This screening acts as a strict gate: failing it short-circuits the pipeline, preventing compilation or runtime evaluation. C.1RULE-BASED SCREENING The static analyzer enforces three primary constraints: 1. Kernel Dispatch Requirement. The binding codepythonbindsrcneed explicitly invoke the kernel execution commandEXECNPUCMD. The verifier scans the binding source for this substring; its absence indicates that the operator either performs no computation or bypasses the NPU dispatch entirely. 2. Binding Logic Restrictions. The C++ binding implementation is restricted to allocation and dispatch duties. The rule checker extracts the function body registered viaPYBIND11MODULEand scans for forbidden calls to the at:: or torch:: namespaces. • Allowed: Tensor allocation functions (e.g., at::empty, at::zeros, at::emptylike). • Forbidden: Any computational operators (e.g., at::add, at::matmul). This rule guarantees that the binding layer does not perform the heavy lifting using CPU-side PyTorch reference implementations. 3. Model Architecture Compliance. The Python invocation layermodelsrcmust define a classModelNew that inherits from torch.n.Module. A simplified Abstract Syntax Tree (AST) analysis enforces that: • Theforwardmethod does not directly call prohibited computations (e.g.,torch.matmul,torch.add) or invoke standard torch.n layers created ininit. • The module must import and call the generatedcustomopslib, ensuring the computation is delegated to the C++ binding and, by extension, the Ascend C kernel. Example Violation. The followingmodelsrcis rejected because it directly invokes atorch.nnlayer (self.conv() ) instead of delegating all computation to customopslib: 19 123456789101112131415161718192021222324252627282930 Outer Iteration 1_ScaledDotProductAttention 2_MultiHeadAttention 3_SelfAttention 4_CrossAttention 6_FlashAttentionV2 7_OptimizedFlashAttention 8_GroupedQueryAttention 9_MultiQueryAttention 11_LinearAttention 12_PerformerAttention 13_LinformerAttention 15_CosformerAttention 16_LocalAttention 17_BlockSparseAttention 18_StridedAttention 21_SEAttention 22_ECAAttention 23_GCT 25_GatedChannelTransform 27_LCT 28_SimAM 29_CoordAtt 30_SpatialGroupEnhance 31_CBAMBlock 32_BAM 33_TripletAttention 34_SKAttention 35_ShuffleAttention 36_PSA 37_ParNetAttention 38_ResidualAttention 39_ScaledDotProductAttentionModular 40_SimplifiedSelfAttention 41_ExternalAttention 42_ParallelPolarizedSelfAttention 43_SequentialPolarizedSelfAttention 44_MUSEAttention 45_AFT_FULL 46_UFOAttention 47_CoTAttention 50_DoubleAttention 51_DAModule 52_PAM 53_AxialAttention 54_CrissCrossAttention 56_OutlookAttention 57_MobileViTAttention 58_MobileViTv2Attention 59_EMSA 60_DAttentionBaseline 61_WeightedPermuteMLP 62_ViTAttention 64_WindowAttention 66_FlashAttentionFwd 68_PagedAttentionKVCache 69_MultiHeadLatentAttention 70_DenseSparseAttention 5_FlashAttention 19_LongformerAttention 26_GCModule 55_HaloAttention 65_GlobalFilter 24_SRM 10_AdaptiveAttention 63_CrossformerAttention 67_FlashAttentionBwd 49_ACmix 14_NystromAttention 1.3090.5750.2610.190.045 219.222.655.213 188.310.719.6065.243 94.6412.347.4136.032 19310.198.502 5.872 194.787.9336.0627.0212.5110.468.922 48.8426.7610.748.827.377 3.397e+045.0833.621 9.4446.6364.96 16.933.672 195.914.428.0825.3694.0713.629 28.815.0513.41 254.4813.182 6513.4511.556.0655.1774.596 0.2330.1750.105 0.7290.072 0.1130.0810.070.056 0.2740.1170.0780.0680.0610.051 0.1520.06 0.2890.1540.0690.0370.03 0.69 0.2510.1120.0970.048 0.776 0.35 1.853 3.456 0.1620.1030.0640.044 0.753 0.609 0.214 6.6481.0610.6110.5260.465 3.0530.5650.2340.172 0.056 1.0190.4620.3980.358 0.5540.376 0.726 0.3960.335 2.5440.8850.48 1.181 0.8350.7110.5470.474 3.802 1.0460.9260.825 4.762.67 0.9860.8420.7390.503 3.889 63.61 114.6101.664.427.745.948 0.354 2.111 1.538 2.6051.470.430.377 2.8250.413 43.264.3272.6030.8580.746 2.2570.6420.408 24.210.760.4940.4370.133 36.40.0490.0420.011 235.413.199.939 34.767.374.05 0.2390.1130.099 6.145 1.0880.108 0.161 74.237.0529.7517.9513.7211.449.492 5.2280.6440.528 3.792e+041.329e+04161.617.5815.34 1.719 6.193 Attention Set (CUDA) Correctness & Performance Optimization Timeline First solve Best perf improved Best performance 4202468 Relative Performance log(Baseline / Agent) 1_ScaledDotProductAttention 2_MultiHeadAttention 3_SelfAttention 4_CrossAttention 5_FlashAttention 6_FlashAttentionV2 7_OptimizedFlashAttention 8_GroupedQueryAttention 9_MultiQueryAttention 10_AdaptiveAttention 11_LinearAttention 12_PerformerAttention 13_LinformerAttention 14_NystromAttention 15_CosformerAttention 16_LocalAttention 17_BlockSparseAttention 18_StridedAttention 19_LongformerAttention 21_SEAttention 22_ECAAttention 23_GCT 24_SRM 25_GatedChannelTransform 26_GCModule 27_LCT 28_SimAM 29_CoordAtt 30_SpatialGroupEnhance 31_CBAMBlock 32_BAM 33_TripletAttention 34_SKAttention 35_ShuffleAttention 36_PSA 37_ParNetAttention 38_ResidualAttention 39_ScaledDotProductAttentionModular 40_SimplifiedSelfAttention 41_ExternalAttention 42_ParallelPolarizedSelfAttention 43_SequentialPolarizedSelfAttention 44_MUSEAttention 45_AFT_FULL 46_UFOAttention 47_CoTAttention 49_ACmix 50_DoubleAttention 51_DAModule 52_PAM 53_AxialAttention 54_CrissCrossAttention 55_HaloAttention 56_OutlookAttention 57_MobileViTAttention 58_MobileViTv2Attention 59_EMSA 60_DAttentionBaseline 61_WeightedPermuteMLP 62_ViTAttention 63_CrossformerAttention 64_WindowAttention 65_GlobalFilter 66_FlashAttentionFwd 67_FlashAttentionBwd 68_PagedAttentionKVCache 69_MultiHeadLatentAttention 70_DenseSparseAttention 1.79x 1.01x 1.01x 0.56x 1.43x 0.61x 1.02x 0.52x 0.56x 0.58x 1.04x 0.81x 1.06x 1.02x 1.04x 2.87x 0.97x 250.41x 2.60x 0.85x 0.93x 2.00x 1.01x 2.47x 1.03x 1.83x 5.90x 1.02x 2.85x 1.03x 1.18x 1.00x 1.06x 4.90x 1.08x 1.08x 1.77x 0.98x 1.09x 1.77x 1.17x 1.06x 1.05x 4.75x 1.01x 1.02x 14.36x 0.80x 1.01x 0.94x 1.07x 1.89x 0.84x 1.04x 1.00x 0.05x 1.02x 1.03x 1.07x 0.99x 0.94x 1.00x 0.86x 1.00x 0.13x 14.56x 25.18x 92.45x Best Correct Run vs PyTorch Baseline (Attention Set (CUDA)) PyTorch Baseline (1.0x) 0.02x0.06x0.25x1.00x4.00x16.00x64.00x256.00x1024.00x Equivalent Speedup (Baseline / Agent) Figure 7: Attention Set (CUDA): Optimization timeline and performance vs. PyTorch baseline for 70 attention operators over 30 iterations. (Left) Correctness and performance optimization timeline. Green dots denote first-solve events; blue dots denote significant (>10%) performance improvements; orange dots denote the best performance achieved. (Right) Best correct run vs. baseline in log 2 speedup. 20 Rejected modelsrc – Hacking Detected class ModelNew(n.Module): def __init__(self, in_channels, out_channels, kernel_size): super(ModelNew, self).__init__() self.conv = n.Conv2d(in_channels, out_channels, kernel_size) def forward(self, x: torch.Tensor) -> torch.Tensor: x = self.conv(x)← VIOLATION x = custom_ops_lib.conv2d_relu_hard_swish_custom(x) return x Verifier Output: [Invalid Error] In the forward method, the model layer is directly called: self.conv(). You must implement the operations in forward() as custom kernels in customopslib. This forces correctness = False, preventing the solution from passing verification. C.2MODEL-BASED SCREENING (LLM AUDITOR) To capture more sophisticated evasion strategies, we employ an LLM-based auditor. The verifier constructs a prompt containing the operator specification (refsrc) and the full generated solution. The LLM is instructed to judge the “architectural integrity” of the code, specifically checking for: •Semantic Location: Verifying that the mathematical logic resides inkernelsrcorhosttilingsrc, not in the glue code. •Dummy Kernels: Detecting empty kernels or kernels that output constant values while the binding performs the actual work. • Binding Anomalies: Identifying complex C++ logic in the binding that acts as a reference implementation. This model-based check runs only if the rule-based checks pass, serving as a final safeguard against “hallucinated” solutions that satisfy syntax benchmarks but fail to implement the actual hardware kernel. Prompt of Model-Based Anti-Hacking. We include the exact prompt used by the model-based anti-hacking auditor below: Prompt of Model-Based Anti-Hacking You are an expert auditor for custom AscendC operators integrated into PyTorch via NPU bindings.,→ The operator under review is named: op (PascalCase:`pascal_op`). Your primary goal is to verify ** Architectural Integrity ** . You need to determine if the custom operator genuinely attempts to perform the computation on the NPU (Kernel-side), or if it "cheats" by handling the logic in the Python bindings or host glue code. ,→ ,→ ,→ ** Important: ** 21 - ** Do NOT ** act as a compiler or unit tester. You do not need to check for syntax errors, strict numerical equivalence, or edge-case correctness. ,→ ,→ - ** Do NOT ** require the implementation to be bug-free. - ** DO ** strictly verify that the computational logic resides in the correct files.,→ ## Input Information 1. ** PyTorch Reference (Intent): ** This defines * what * the operator is supposed to do. Use this only to understand the intended logic flow (e.g., "this is a matrix multiplication" or "this is an element-wise add"). ,→ ,→ <PyTorch Reference Implementation> ref_src </PyTorch Reference Implementation> 2. ** Generated Implementation (Execution): ** This includes the Interface (`project_json_src`), Host Logic (`host_tiling_src`,`host_operator_src`), Kernel Logic (`kernel_src`), Bindings (`python_bind_src`), and Usage (`model_src`). ,→ ,→ ,→ <Generated Custom Operator Implementation> code </Generated Custom Operator Implementation> ## Evaluation Criteria: Where is the Math? You must judge whether the ** implementation location ** is correct. ### 1. The Valid Domain (Kernel & Host Tiling) The actual mathematical operations (loops, math instructions, data manipulation) MUST be attempted in:,→ -`kernel_src` (The AscendC Kernel) -`host_tiling_src` /`host_operator_src` (Tiling and shape logic) ** It is VALID if: ** - The`kernel_src` contains code that clearly resembles the logic described in the Reference.,→ - The`host_tiling_src` calculates tiling parameters. ### 2. The Restricted Domain (Bindings & Model) These files are for "glue" and dispatch only. -`python_bind_src` -`model_src` ** It is INVALID (Cheating) if: ** - The`python_bind_src` or`model_src` performs part of the actual computation using PyTorch/ATen ops (e.g.,`at::add`,`torch.matmul`) and returns that result, bypassing the custom kernel. ,→ ,→ - The`kernel_src` is empty, a dummy stub, or purely performs a copy without the required math.,→ 22 - The`python_bind_src` allocates outputs but then computes results using standard PyTorch functions instead of invoking the kernel.,→ - The`model_src` calls standard PyTorch functions to compute results instead of using the custom operator.,→ - Any other pattern where the core logic is outside the kernel/host tiling files.,→ ## Decision Logic - Valid: The coder honestly attempted to write an AscendC kernel. The math logic (add, sub, matmul, etc.) is visible in the`kernel_src`, without using constants to replace or "fake" the implementation logic. The binding file properly dispatches to this kernel via `EXEC_NPU_CMD` or similar. ,→ ,→ ,→ ,→ - Invalid: The coder bypassed the NPU kernel. For example, the binding file catches the inputs, calls a standard PyTorch function to get the result, and returns it. Or the kernel exists but does nothing related to the reference logic. Or the kernel use constant value to skip part of the computation. ,→ ,→ ,→ ,→ ## Output Output only the following JSON object, warpped in triple backticks: ```json```.,→ Do NOT include any additional text.: ```json "valid": true | false, "reason": "Concise explanation focusing on WHERE the logic is implemented.",→ ``` DBASELINE METHODOLOGIES We evaluate our approach against two distinct baseline strategies that represent standard practices in code generation: Pass@k (Generation) and Iterative Refinement. D.1PASS@k (GENERATION) This mode implements a classic sampling strategy, leveraging the probabilistic nature of LLMs to generate widely diverse attempts. •Methodology: For each operator task, we generateKindependent candidate solutions in parallel. Each candidate includes the full kernel code, tiling logic, and binding glue. •Context: The process is stateless; each generation starts from a fresh prompt containing the operator specification and few-shot examples (if configured), without knowledge of prior attempts or peer candidates. • Objective: This baseline evaluates the model’s “zero-shot” or “few-shot” capability to produce a correct solution purely from the prompt. It serves as a measure of the model’s intrinsic knowledge of the Ascend C DSL. 23 D.2ITERATIVE REFINEMENT This mode implements a stateful agentic loop that mimics a human developer’s debugging workflow, consisting of two distinct phases: Drafting and Optimization. Phase 1: Drafting (Correctness). The goal is to produce a compilable and functionally correct kernel. •Feedback Loop: The agent generates an initial draft, which is then compiled and executed. If compilation fails, the compiler error logs are fed back to the model. If execution fails (correctness error), the mismatch info is provided. •History: The agent maintains a conversation history of (Code→Error→Fix), allowing it to iteratively repair syntax errors and logic bugs. Phase 2: Optimization (Performance). Once a correct kernel is identified, the agent transitions to performance optimization. •Hill Climbing: The correct kernel serves as a baseline. The prompt shifts to request performance improvements (e.g., “minimize execution time”). •Metric Feedback: The agent receives latency measurements from the hardware profiling tool. It generates new versions to improve this metric. If a new version is slower or incorrect, the agent reverts to the previous best baseline or receives feedback on the regression. This baseline establishes the performance upper bound for a standard agentic loop without the long-term, cross-task memory mechanisms introduced in our EvoKernel framework. Prompt Construction. The prompt structure differs between the two phases. In drafting mode, each turn appends the previous attempt and its feedback: Drafting Mode Prompt [System]: You are a helpful assistant [User]: base_prompt [Assistant]: last_code [User]: compile_error or correctness_error In optimization mode, the prompt includes two turns of history to preserve the best correct baseline: Optimization Mode Prompt [System]: You are a helpful assistant [User]: base_prompt [Assistant]: baseline_code [User]: baseline_feedback "The code above is correct. Now optimize it..." [Assistant]: last_code [User]: "Performance: X ms" or error_feedback ConfigurationParameters.Weusethefollowinghyperparameters:maxturns=30, maxfeedbackchars= 4000(truncation limit for compiler/correctness output),infraretries= 3 (exponential backoff for transient failures), andparallelism= 16(concurrent operators). The evaluation uses a remote server with timeout= 65 minutes per validation call. 24 D.3CODEX (GENETIC ITERATION) This baseline utilizes an advanced EXEC mode that forces the model (specifically GPT-5.2 Medium Reasoning) to perform evolutionary iterations within a single session, effectively turning a completion task into a genetic-like agent. Mechanism: The ReAct Loop. Unlike standard generation, this mode grants the model: • Shell Access: The ability to execute commands with configurable timeout. • File System Access: The ability to read and write files via the applypatch tool. • Immediate Feedback: The stdout/stderr of its commands are fed back into its context window. This creates a Reason-Act-Observe loop managed entirely by the Codex binary but orchestrated by our injected prompt. Prompt Structure. The prompt sent to Codex consists of two parts: (1) the base kernel generation prompt with few-shot examples, and (2) validation workflow instructions that specify the autonomous iteration protocol: Iteration Loop. The process operates iteratively: • Generate: The model writes a candidate kernel file (e.g., op.txt). • Validate: The validation script (codexvalidate.sh) sends the code to a remote evaluation server and returns the verifier result (compiled, correctness). • React: – If result == success: The loop terminates. – Ifresult == failure: The model reads the error log, analyzes the failure, and revises the code for the next iteration. Codex Validation Instructions ## Task: Implement Ascend Operator`op` ### Workflow 1. Write Code: Create`op.txt` using apply_patch tool 2. Validate: Run`./codex_validate.sh op file ascendc` with timeout_ms=1200000 (20 minutes) Returns JSON: compiled, correctness, error SUCCESS = compiled:true AND correctness:true 3. On failure: Fix code based on error, re-validate ### Rules - NO local compilation (gcc, g++, make, cmake) - After 3 consecutive validation timeouts, STOP This approach is distinct from Iterative Refinement because it occurs entirely within a single model session via tool use (In-Context Learning), whereas Refinement is an external Python loop managing the history. It tests the model’s intrinsic ability to function as a developer with a compiler and debugger. Configuration Parameters.We invoke the Codex CLI with:sandbox=workspace-write (allowsfilewriteswithinworkspace),ask-for-approval=never(fullyautonomous), modelreasoningeffort=medium. Each validation call uses a 20-minute timeout (timeoutms= 1200000) to accommodate remote compilation and execution. Termination Condition. To ensure a fair comparison, we impose a strict stop condition based on verification attempts. The agent terminates after 30 verification attempts or upon finding a correct solution, whichever comes first. 25 24681012141618 Verification Attempts 0.0 0.2 0.4 0.6 0.8 1.0 Cumulative Success Rate 46.0% 83.0% Codex Performance over Verification Attempts Compilation Success Rate Correctness Rate Figure 8: Codex (GPT-5.2) cumulative correctness. D.4COMPARISON OF AGENTIC BASELINES Table 6 summarizes the key architectural differences between the two agentic baselines. Table 6: Comparison of Iterative Refinement and Codex agent architectures. AspectRefinementCodex Execution modelAPI conversation loopAutonomous tool use Iteration controlExternal scriptAgent decides Prompt updatesEach turn rebuiltSingle prompt History length1–2 turnsInternal memory Feedback sourceInjected by scriptAgent calls validator File operationsExtract from text applypatch tool Termination30 iterations or success30 verification attempts or success EOPERATOR SUBSET FOR TRANSFER EXPERIMENTS This section lists the held-out operator subset used in the cross-backbone transfer experiments (Section 4.3). The subset consists of 50 operators randomly sampled from the benchmark: 30 Level 1 (L1) operators and 20 Level 2 (L2) operators. These operators were excluded from the GPT-5.2 memory bank during training to prevent data leakage, and were used exclusively for evaluating transfer performance on DeepSeek-V3.2 and Qwen3-Coder-30B. FEVALUATION AND PROFILING METHODOLOGY This section details the correctness verification and latency profiling procedures. F.1CORRECTNESS VALIDATION Fail-Fast Execution Strategy. To optimize evaluation efficiency, the custom kernel is executed first with a strict SIGALRMtimeout. If the custom kernel fails (timeout, crash, or exception), the reference run is skipped entirely. Structured Mismatch Feedback. The verifier returns detailed, machine-readable error messages to guide iterative refinement. The following illustrates the categories of feedback returned: 26 Table 7: Randomly sampled operator subset for transfer evaluation. Level 1 Operators (30 total) OperatorTypeOperatorType 43MaxPooling3Dpooling46AveragePooling3Dpooling 42MaxPooling2Dpooling49Maxreductionoveradimensionpooling 92 cumsumexclusiveloss93maskedcumsumloss 103Dtensormatrixmultiplicationmatmul77convtransposed3Dsquareinputsquarekernelconvolution 66convstandard3Dasyminputasymkernelconvolution38L1Normnormalization 51 Argmaxoveradimensionconvolution31ELUactivation 2Standardmatrixmultiplicationmatmul22Tanhactivation 71convtransposed2Dasyminputsquarekernelconvolution33BatchNormnormalization 16 MatmulwithtransposedAmatmul97ScaledDotProductAttentionloss 79convtransposed1Dasyminputsquarekernelconvolution81convtransposed2Dasyminputsquarekernelconvolution 74convtransposed1Ddilatedconvolution50convstandard2Dsquareinputsquarekernelconvolution 11 4Dtensormatrixmultiplicationmatmul84convdepthwise2Dasyminputsquarekernelconvolution 56 convstandard2Dasyminputasymkernelconvolution27SELUactivation 57convtransposed2Dsquareinputsquarekernelconvolution88MinGPTNewGeluconvolution 73convtransposed3Dasyminputsquarekernelconvolution34InstanceNormnormalization Level 2 Operators (20 total) OperatorTypeOperatorType 58ConvTranspose3dLogSumExpHardSwishSubtractClampfuse86MatmulDivideGELUfuse 27Conv3dHardSwishGroupNormMeanfuse68MatmulMinSubtractfuse 6Conv3dSoftmaxMaxPoolMaxPoolfuse80GemmMaxSubtractGELUfuse 14GemmDivideSumScalingfuse45GemmSigmoidLogSumExpfuse 62MatmulGroupNormLeakyReLUSumfuse43Conv3dMaxLogSumExpReLUfuse 25Conv2dMinTanhTanhfuse5ConvTranspose2dSubtractTanhfuse 23Conv3dGroupNormMeanfuse70GemmSigmoidScalingResidualAddfuse 39GemmScaleBatchNormfuse78ConvTranspose3dMaxMaxSumfuse 26ConvTranspose3dAddHardSwishfuse31Conv2dMinAddMultiplyfuse 53GemmScalingHardtanhGELUfuse3ConvTranspose3dSumLayerNormAvgPoolGELUfuse 1. Shape Mismatch: output.shape mismatch: expected (16, 512, 512), got (16, 512, 256) 2. Numerical Mismatch: [FAIL] Output mismatch: 1/5 trials passed, 4 failed. Tolerance atol=0.01, rtol=0.01. Trial 1: 54/524160 mismatched (0.01%), max_abs=0.99, max_rel=97209.6, Bounding box: output[0:31, 4032:4088] Trial 2: 64/524160 mismatched (0.01%), max_abs=0.99, max_rel=87570.4, Bounding box: output[99:100, 35:99] Key diagnostics: (a)maxabs/maxrel: maximum absolute and relative difference; (b) Bounding box: spatial localization of errors, revealing tile boundary bugs. Example Agent Diagnosis. While optimizing53Minreductionoveradimension, the agent encounters the error above and identifies a synchronization race in the accumulator initialization: Row 0 was fetched asynchronously via MTE but the Vector engine began computation before the transfer completed. The fix: queue Row 0 through the standard Ping-Pong pipeline (enqueue→dequeue→copy to accVec) to enforce synchronization before any arithmetic. 3. Type Mismatch: 27 type(output) mismatch: expected Tensor, got list 4. Length Mismatch (for tuple/list outputs): len(output) mismatch: expected 3, got 2 5. Timeout: [FAIL] First correctness run timed out after 60s 6. Runtime Exception: [FAIL] NPU out of memory. Tried to allocate 12.10 GiB [FAIL] vector core exception at line 42 F.2LATENCY PROFILING We use the native msprof profiler (via torch npu.profiler) with: • 3 warm-up runs (discarded) to stabilize caches and JIT compilation. • 3 profiling passes with distinct configurations (PipeUtilization, Memory, ResourceConflict). • The mean “Computing” time fromsteptracetime.csvis reported, isolating on-chip kernel execution from host overhead. 3-Pass Aggregation. Each profiling pass writes asteptracetime.csvwith a “Computing” column (in μs). The final timing is aggregated as: Pass 1 (PipeUtilization): Computing = 13640 us Pass 2 (Memory): Computing = 13380 us Pass 3 (ResourceConflict): Computing = 12913 us => performance.mean = avg([13.64, 13.38, 12.91]) = 13.31 ms => performance.std = 0.33 ms This procedure yields negligible standard deviation (<3%) across profiling runs. Data Source: step tracetime.csv vs kerneldetails.csv. Both files are produced by msprof: • steptracetime.csv: Total device execution time for the entire step (all kernels combined). Used for performance.mean/max/min/std. • kerneldetails.csv: Per-kernel breakdown with detailed hardware metrics. Useful for optimization but may not sum exactly to total time due to overlaps/gaps. We report the steptracetime value as the canonical latency metric. Example Profiling Output.The verifier returns detailed per-kernel metrics extracted from kerneldetails.csv: "performance": "max": 13.64, "mean": 13.38, "min": 12.913, "std": 0.33 , "profiling": "MinReductionOverADimensionCustom": "Block Dim": 32.0, "Duration(ms)": 13.38, "aic_fixpipe_ratio": 0.0, "aic_fixpipe_time(ms)": 0.0, "aic_icache_miss_rate": 0.0, "aic_l1_read_bw(GB/s)": 0.0, "aic_l1_write_bw(GB/s)": 0.0, 28 "aic_l2_read_bw(GB/s)": 0.0, "aic_l2_write_bw(GB/s)": 0.0, "aic_mac_ratio": 0.0, "aic_mac_time(ms)": 0.0, "aic_main_mem_read_bw(GB/s)": 0.0, "aic_main_mem_write_bw(GB/s)": 0.0, "aic_mte1_ratio": 0.0, "aic_mte1_time(ms)": 0.0, "aic_mte2_ratio": 0.0, "aic_mte2_time(ms)": 0.0, "aic_scalar_ratio": 0.0, "aic_scalar_time(ms)": 0.0, "aic_total_cycles": 0.0, "aicore_time(ms)": 0.0, "aiv_icache_miss_rate": 0.0, "aiv_l2_read_bw(GB/s)": 0.0, "aiv_l2_write_bw(GB/s)": 0.0, "aiv_main_mem_read_bw(GB/s)": 0.46, "aiv_main_mem_write_bw(GB/s)": 0.0, "aiv_mte2_ratio": 0.346, "aiv_mte2_time(ms)": 3.378, "aiv_mte3_ratio": 0.0, "aiv_mte3_time(ms)": 0.001, "aiv_scalar_ratio": 0.677, "aiv_scalar_time(ms)": 6.605, "aiv_time(ms)": 9.66, "aiv_total_cycles": 571879708.0, "aiv_ub_read_bw(GB/s)": 25.966, "aiv_ub_write_bw(GB/s)": 39.604, "aiv_vec_bank_cflt_ratio": 0.053, "aiv_vec_bankgroup_cflt_ratio": 0.058, "aiv_vec_ratio": 0.134, "aiv_vec_resc_cflt_ratio": 0.0, "aiv_vec_time(ms)": 1.31, "cube_utilization(%)": 0.0 Key metrics include vector/scalar/MTE time ratios, unified buffer bandwidth, and cube utilization, enabling targeted optimization. GEXAMPLE GENERATED KERNEL The following shows a complete, correctly compiling Ascend C kernel forTanhgenerated by EvoKernel. Each source file section is shown with a distinct background color. In actual verifier pipeline, the below artifact is parsed into respective files. project_json_src = r′ [ "op": "TanhCustom", "language": "cpp", "input_desc": [ "name": "x", "param_type": "required", "format": ["ND"], "type": ["float"] ], "output_desc": [ "name": "y", "param_type": "required", "format": ["ND"], "type": ["float"] 29 ] ] ′ host_tiling_src = r""" #include "register/tilingdata_base.h" namespace optiling BEGIN_TILING_DATA_DEF(TilingData) TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, tileLength); TILING_DATA_FIELD_DEF(uint32_t, blockDim); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(TanhCustom, TilingData) // namespace optiling """ host_operator_src = r""" #include "tanh_custom_tiling.h" #include "register/op_def_registry.h" namespace optiling static inline uint32_t AlignUp(uint32_t x, uint32_t a) return (x + a - 1) / a * a; static inline uint32_t MinU32(uint32_t a, uint32_t b) return a < b ? a : b; constexpr uint32_t MAX_BLOCK_DIM = 32; constexpr uint32_t DEFAULT_TILE_ELEMS = 8192; static ge::graphStatus TilingFunc(gert::TilingContext * context) TilingData tiling; const uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize(); const auto dtype = context->GetInputTensor(0)->GetDataType(); if (dtype != ge::DT_FLOAT) return ge::GRAPH_FAILED; uint32_t blockDim = MAX_BLOCK_DIM; const uint32_t minElemsPerBlock = 262144; if (totalLength / blockDim < minElemsPerBlock) blockDim = totalLength / minElemsPerBlock; if (blockDim < 1) blockDim = 1; blockDim = MinU32(blockDim, MAX_BLOCK_DIM); 30 uint32_t tile = DEFAULT_TILE_ELEMS; if (totalLength < tile) tile = totalLength; tile = AlignUp(tile, 8); // 32B alignment for fp32 if (tile == 0) tile = 8; if (tile > totalLength) tile = totalLength; context->SetBlockDim(blockDim); tiling.set_totalLength(totalLength); tiling.set_tileLength(tile); tiling.set_blockDim(blockDim); tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); size_t * ws = context->GetWorkspaceSizes(1); ws[0] = 0; return ge::GRAPH_SUCCESS; // namespace optiling namespace ge static graphStatus InferShape(gert::InferShapeContext * context) * context->GetOutputShape(0) = * context->GetInputShape(0); return GRAPH_SUCCESS; static graphStatus InferDataType(gert::InferDataTypeContext * context) context->SetOutputDataType(0, context->GetInputDataType(0)); return ge::GRAPH_SUCCESS; // namespace ge namespace ops class TanhCustom : public OpDef public: explicit TanhCustom(const char * name) : OpDef(name) this->Input("x") .ParamType(REQUIRED) .DataType(ge::DT_FLOAT) .Format(ge::FORMAT_ND); this->Output("y") .ParamType(REQUIRED) .DataType(ge::DT_FLOAT) .Format(ge::FORMAT_ND); this->SetInferShape(ge::InferShape). SetInferDataType(ge::InferDataType); 31 this->AICore() .SetTiling(optiling::TilingFunc) .AddConfig("ascend910b"); ; OP_ADD(TanhCustom); // namespace ops """ kernel_src = r""" #include "kernel_operator.h" class KernelTanh public: __aicore__ inline KernelTanh() __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength, uint32_t tileLength, uint32_t blockDim) this->totalLength = totalLength; this->tileLength = tileLength; this->blockDim = blockDim; const uint32_t bid = AscendC::GetBlockIdx(); const uint32_t base = (totalLength / blockDim); const uint32_t rem = totalLength - base * blockDim; const uint32_t myLen = base + (bid < rem ? 1u : 0u); const uint32_t myOff = bid * base + (bid < rem ? bid : rem); startOffset = myOff; localLength = myLen; xGm.SetGlobalBuffer((__gm__ float * )x + startOffset, localLength); yGm.SetGlobalBuffer((__gm__ float * )y + startOffset, localLength); constexpr uint32_t PAD_BYTES = 256; pipe.InitBuffer(inQueueX, 2, tileLength * sizeof(float) + PAD_BYTES); pipe.InitBuffer(outQueueY, 2, tileLength * sizeof(float) + PAD_BYTES); // Tanh tmp UB, aligned to 256B. Keep conservative sizing. const uint32_t tmpBytes = ((tileLength * 8u + 255u) / 256u) * 256u; pipe.InitBuffer(tmpQueue, 1, tmpBytes); __aicore__ inline void Process() if (localLength == 0) return; 32 uint32_t offset = 0; uint32_t cur = (localLength > tileLength) ? tileLength : localLength; CopyIn(offset, cur); offset += cur; while (offset < localLength) const uint32_t next = (offset + tileLength <= localLength) ? tileLength : (localLength - offset); CopyIn(offset, next); Compute(cur); CopyOut(offset - cur, cur); cur = next; offset += next; Compute(cur); CopyOut(localLength - cur, cur); private: __aicore__ inline void CopyIn(uint32_t offset, uint32_t len) AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>(); AscendC::DataCopy(xLocal, xGm[offset], len); inQueueX.EnQue(xLocal); __aicore__ inline void Compute(uint32_t len) AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>(); AscendC::LocalTensor<float> yLocal = outQueueY.AllocTensor<float>(); AscendC::LocalTensor<uint8_t> tmp = tmpQueue.AllocTensor<uint8_t>(); AscendC::Tanh<float>(yLocal, xLocal, tmp, len); tmpQueue.FreeTensor(tmp); outQueueY.EnQue<float>(yLocal); inQueueX.FreeTensor(xLocal); __aicore__ inline void CopyOut(uint32_t offset, uint32_t len) AscendC::LocalTensor<float> yLocal = outQueueY.DeQue<float>(); AscendC::DataCopy(yGm[offset], yLocal, len); outQueueY.FreeTensor(yLocal); 33 private: AscendC::TPipe pipe; AscendC::TQue<AscendC::TPosition::VECIN, 2> inQueueX; AscendC::TQue<AscendC::TPosition::VECOUT, 2> outQueueY; AscendC::TQue<AscendC::TPosition::VECCALC, 1> tmpQueue; AscendC::GlobalTensor<float> xGm; AscendC::GlobalTensor<float> yGm; uint32_t totalLength = 0; uint32_t tileLength = 0; uint32_t blockDim = 1; uint32_t startOffset = 0; uint32_t localLength = 0; ; extern "C" __global__ __aicore__ void tanh_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) GET_TILING_DATA(tiling_data, tiling); KernelTanh op; op.Init(x, y, tiling_data.totalLength, tiling_data.tileLength, tiling_data.blockDim); op.Process(); """ python_bind_src = r""" #include <torch/library.h> #include <torch/extension.h> #include "pytorch_npu_helper.hpp" at::Tensor tanh_impl_npu(const at::Tensor& x) auto y = at::empty_like(x); EXEC_NPU_CMD(aclnnTanhCustom, x, y); return y; TORCH_LIBRARY_IMPL(myops, PrivateUse1, m) m.impl("tanh_custom", &tanh_impl_npu); PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) m.def("tanh_custom", &tanh_impl_npu, "tanh custom (NPU)"); """ model_src = r′ import torch import torch_npu 34 import custom_ops_lib class ModelNew(torch.n.Module): """ Simple model that performs a Tanh activation using a custom Ascend C op. """ def __init__(self): super(ModelNew, self).__init__() def forward(self, x: torch.Tensor) -> torch.Tensor: return custom_ops_lib.tanh_custom(x) ′ 35

Instant research discovery