Source-linked AI summary
Beyond Scaling: Self-Evolving LLM Agents for Hardware Kernel Optimization via an Experience-Driven Workflow and Experience Graph Memory
Siyuan Chen, Runlin Hou, Shenxiu Wu, Yansong Sun, Junming Cao, Yiyu Zhang, Shudi Shao, Junhao Qiu, Zhichao Lu, Qingfu Zhang
TL;DR
Hardware-kernel agents rarely turn execution feedback into reusable evidence, especially when target hardware has sparse public optimization data. KOPE records decisions and outcomes in Experience Graph Memory and injects relevant experience under a fixed token budget, achieving stronger reported coverage and performance while the foundation model remains fixed.
Problem
Existing kernel-optimization agents rarely organize execution feedback as reusable evidence for later decisions or tasks, while sparse target-hardware data limits available model knowledge.
Method
KOPE combines Experience Graph Memory for ordered decision-to-outcome histories and alternative branches with Active Context Management and Injection for bounded, task-relevant retrieval.
Results
KOPE achieves higher full-suite pass rates with both evaluated models and a 1.54× per-operator speedup geometric mean relative to CANNBot with GLM-5.2.
Takeaways & Limitations
The results support continual hardware-kernel optimization through external experience while the foundation model remains fixed.
Takeaways & Limitations
The system comparison changes several workflow decisions at once, so its gain is not attributed to any individual component.
Abstract
from arXiv · showhide
Hardware kernel optimization requires repeated compilation, correctness testing, profiling, and revision. LLM agents can automate parts of this process, and stronger foundation models, longer context windows, and longer execution horizons have improved optimization within individual tasks. These advances alone do not enable an agent to learn from completed optimization runs. Existing kernel-optimization agents seldom preserve a decision, its observed execution feedback, and the later decisions that use that evidence. Retaining every prior trajectory is also impractical because an expanding history competes with the current task for context. We present KOPE, an experience-driven framework for hardware kernel optimization. KOPE records optimization trajectories with correctness and performance feedback in Experience Graph Memory, then uses Active Context Management and Injection to retrieve relevant experience under a fixed token budget. The graph retains decision order, observed outcomes, and alternative branches, allowing evidence collected on the target hardware to inform later optimization steps and tasks. Under the same GLM-5.2 setting, the geometric mean of KOPE's per-operator speedups is $1.54\times$ that of CANNBot, the strongest competing baseline. In a complete 53-operator ablation, Active Context Management and Injection raises pass rate from 60.0\% to 84.6\%, increases the evaluator-reported positive-field geometric mean from 0.0382 to 0.0661, and reduces optimization token consumption from 15.9B to 1.113B tokens relative to passive agent-led context construction. Enabling Experience Graph Memory raises full-suite pass rate from 55.2\% to 84.6\% and yields a $1.43\times$ geometric-mean speedup on valid timing comparisons. These results support continual optimization through external experience while the foundation model remains fixed.
Introduction
KOPE addresses the limited reuse of execution feedback in hardware-kernel agents by retaining optimization evidence externally and injecting relevant experience into later decisions. Its workflow targets hardware with sparse public optimization data and reports stronger coverage and efficiency than the cited baselines and ablations.
- Motivation: Existing agents mainly use execution feedback within the current run, leaving later decisions dependent on knowledge already encoded in the model.Expanding the retained trajectory can also consume context and obscure relevant evidence.
- Approach: KOPE converts compiler diagnostics, correctness outcomes, profiling observations, and performance changes into external knowledge while keeping model parameters fixed.Experience Graph Memory preserves decision-to-outcome histories and alternative branches.
- Approach: The workflow validates candidates for correctness and performance, records trajectories, and reuses the resulting evidence in later decisions and runs.The evaluation uses Huawei Ascend NPUs and AscendC kernels, where public implementations and optimization traces are limited.
- Results: 1.54× is KOPE’s per-operator speedup geometric mean relative to CANNBot with GLM-5.2.The complete 53-operator ablation also reports pass rate rising from 60.0% to 84.6% and token consumption falling from 15.9B to 1.113B.
Related Work
Prior work strengthens kernel optimization through model capability, adaptation, agentic workflows, retrieval, memory, and context management. These lines of work motivate KOPE’s focus on reusable execution evidence and effective use of bounded context.
- Evaluation setting: CANN Bench extends execution-based evaluation to Ascend NPUs with 53 operators, 1,060 public test cases, and a composite compilation, correctness, and performance score.This benchmark provides the evaluation setting for the paper’s Ascend-focused study.
- Evaluation setting: The appendix reports additional Triton-kernel optimization results on RISC-V hardware to assess generalization beyond the primary AscendC environment.The passage identifies this as an assessment beyond the primary hardware and software environment.
- LLM-based kernel optimization: Kernel-optimization research combines domain adaptation, reinforcement learning, instruction tuning, multi-agent planning, and execution-based evaluation.Examples include Astra, Makora, CuSeT, QiMeng-Kernel, and AscendKernelGen.
- Retrieval and agent memory: Retrieval and agent-memory systems pair generation with external information, compact prompts, persistent interactions, or distilled task trajectories.The cited examples include RAG, RepoCoder, hierarchical context pruning, MemGPT, and ExpeL.
- Long-context utilization: Longer context windows do not guarantee effective use of all included information, motivating selective context construction.“Lost in the Middle” and RULER report degraded or shorter-than-advertised usable context on several tasks.
KOPE: Experience-Driven Agent Workflow
KOPE forms a closed optimization loop in which target-side execution feedback updates persistent experience, and graph-aware retrieval injects bounded, state-relevant evidence into each subsequent decision. Its memory preserves outcomes, provenance, and branching while context assembly enforces feasibility and token limits.
- Learning from Exploration: Each optimization attempt maps the current task state and persistent memory to retrieved experience, a bounded prompt, an agent action, execution outcomes, and updated memory.The fixed-parameter agent receives compiler diagnostics, correctness results, and measured performance from execution.
- Learning from Exploration: KOPE compiles and tests every candidate, replacing the incumbent only when a correct candidate improves the measured objective.Successful transformations become positive evidence, while failures and regressions remain available as failure evidence.
- Experience Graph Memory: Experience Graph Memory records decisions with observed outcomes and provenance, connecting execution feedback to later decisions within and across tasks.The graph stores alternative branches and preserves execution order without asserting causal identification.
- Experience Graph Memory: KOPE-Mem stores each event as both an append-only Journal and a structured Case containing strategy, validation, measurements, feedback, confidence, and lineage identifiers.The dual representation supports narrative summaries alongside structured retrieval fields.
- Experience Graph Memory: The reconstructed graph is a directed acyclic forest in which each Case has at most one predecessor and may have several successors.Alternative attempts can share a predecessor, while edges record lineage rather than causal effects.
- Graph-aware retrieval and ranking: Graph-aware retrieval combines exact filters and semantic search with a separate downstream-outcome signal for ranking relevant Cases.The outcome label is an observed step-outcome category, and correctness remains a separate field.
- Active Context Management and Injection: KOPE rebuilds the prompt before every iteration so the query reflects the incumbent, latest diagnostics, and current optimization target.An event recorded at iteration k can alter the context supplied at iteration k + 1.
- Dynamic context tiers: Hot context holds required task state, warm context holds retrieved Cases and summaries, and cold context holds broader documentation.The tiers define admission priority for the current prompt rather than permanent storage locations.
Experimental Setup
The evaluation fixes the AscendC benchmark, hardware, models, and token budget, then compares KOPE with baselines and tests its context and graph-memory components using full-suite artifacts.
- Benchmark and Hardware: The benchmark contains 53 AscendC operators and 1,060 public cases evaluated on Ascend 910C hardware.Each operator has 20 cases, and missing operators count as failures.
- Systems and Baselines: KOPE, CANNBot, and CUDA-Agent are compared under fixed models, benchmark, and 1M-token budgets.The comparison uses GLM-5.2 and Deepseek-V4-Pro with the same Ascend evaluation workflow.
- Ablation Design: The active-context ablation compares passive agent-constructed context with active selection, management, and injection of task-relevant experience across all 53 operators.Both configurations use GLM-5.2 and archived artifacts covering the complete operator set.
- Ablation Design: The graph-memory ablation disables or enables Experience Graph Memory while holding the GLM-5.2 workflow and target kernels fixed.Four artifacts span all 53 operators and provide paired implementations for comparison.
- Metrics and Evaluation: Pass rate counts accepted cases over 1,060, while geometric-mean speedup excludes missing and nonpositive fields and is not interpreted as suite-wide coverage.The evaluator checks compilation, correctness, performance, and reward-hacking safeguards before accepting cases.
- Artifacts and Reporting: Table 1 reports the best completed final-evaluation artifact for each model-system configuration, with missing operators counted as failures.The active/passive result uses one complete configuration pair, and the graph-memory timing result uses 412 paired valid comparisons.
Results
KOPE improves full-suite kernel-optimization coverage and performance over competing workflows, while its active context mechanisms and experience graph substantially improve correctness and efficiency. These gains are measured across complete-suite comparisons and ablations, with some speedup estimates conditional on matched positive timings.
- KOPE’s GLM-5.2 operator-speedup geometric mean is 1.54× CANNBot’s under the headline aggregation.
- KOPE passes 897 of 1,060 cases, versus CANNBot’s 613, yielding a 57.8% full-suite pass rate after missing operators count as failures.On the exact 49-operator intersection, KOPE passes 86.3% of cases versus CANNBot’s 62.6%.
- KOPE’s Deepseek-V4-Pro full-suite pass rate is 73.9%, compared with CANNBot’s 39.3%, although CANNBot remains competitive on the returned subset.KOPE provides broader functional coverage and a higher full-suite score, while CANNBot has a slightly higher positive-field geometric mean on the 31-operator intersection.
- Active Context Management and Injection: Active Context Management and Injection raises complete-suite pass rate from 60.0% to 84.6% and positive-field geometric mean from 0.0382 to 0.0661.The benchmark score rises from 636.00 to 2004.49, a 1.73× increase in positive-field geometric mean.
- Active Context Management and Injection: Active context construction reduces aggregate optimization token consumption from 15.9B to 1.113B, a 93.0% reduction, with both artifacts covering all 53 operators.The comparison therefore requires neither failure imputation for missing outputs nor intersection selection.
- Experience Graph Memory: Experience Graph Memory raises full-suite pass rate from 55.2% to 84.6% and yields a 1.434× geometric-mean speedup across 412 timing pairs.The acceleration claim is conditional on the 38.9% of CANN Bench with positive timings in both configurations, and the L4 estimate uses only 28 paired cases.
Discussion
KOPE’s system-level advantage is strongest in complete-suite coverage, but the comparisons do not isolate individual workflow components and do not guarantee superior conditional performance.
- System-level comparison: 26.8 percentage points is KOPE’s GLM-5.2 full-suite pass-rate lead over CANNBot, while Deepseek-V4-Pro shows a 34.5-point lead.On the 31-operator Deepseek-V4-Pro intersection, CANNBot has the higher conditional score and slightly higher positive-field geometric mean.
- Attribution limits: The system comparison changes several workflow decisions simultaneously, so it cannot attribute gains to context assembly, trajectory ranking, or another component.The available evidence supports associations between active context construction, coverage and token use, and between graph memory, pass rate and paired performance.
- System-level comparison: 9.1% to 14.7% is CUDA-Agent’s full-suite pass-rate range, and neither configuration solves a complete operator.Its conditional speedups vary sharply by model, so reached-subset performance does not represent suite-wide coverage.
- Evidence scope: One complete configuration pair underlies the active/passive context result, while the graph-memory result uses one correctness comparison and one 412-case timing comparison.Repeated paired runs are required to estimate run-to-run variance.
Conclusion
KOPE externalizes target-side optimization experience in a graph and selectively reinjects it under a fixed token budget. Its ablations and system comparisons support improved coverage and aggregate performance, while the evidence remains based on archived runs.
- Conclusion: KOPE-Mem records ordered decision-to-outcome histories and alternative branches, while Active Context Management and Injection supplies relevant experience under a fixed token budget.The workflow keeps the foundation model fixed while turning exploration into reusable optimization evidence.
- Conclusion: Active context construction raises full-suite pass rate from 60.0% to 84.6%, improves the positive-field geometric mean by 1.73×, and reduces tokens from 15.9B to 1.113B.This comparison is against passive agent-led context construction.
- Conclusion: Experience Graph Memory raises pass rate from 55.2% to 84.6% and yields a 1.434× geometric-mean speedup on 412 paired timing cases.At system level, KOPE’s GLM-5.2 per-operator speedup geometric mean is 1.54× that of CANNBot.
- Conclusion: The results support retaining and selectively reinjecting target-side experience for hardware with sparse public training data.Each configuration is represented by one archived run.
KOPE on RISC-V
The RISC-V evaluation retargets KOPE to Triton kernels on a SpacemiT K3 processor and compares four knowledge settings using common-case coverage and fixed-intersection performance metrics.
- Experimental setup: KOPE is evaluated on Triton kernels from FlagGems running on a SpacemiT K3 RISC-V processor and compared with PyTorch native kernels.This appendix changes the hardware, kernel language, and software stack from the primary AscendC environment.
- Evidence boundary: One archived result per setting limits the evidence, with one latency value per case and no numerical-error traces retained.The legacy calibration metadata is not read in analysis or aggregation.
- Benchmark: 204 cases from 21 operators form the common universe, covering fp32 and fp16 while disabling bf16 and capping tensors at 16.7M elements.Timing uses triton.testing.do_bench on the kernel-only core suite.
- Aggregation: Performance is reported both as a conditional geometric mean over each setting’s passed cases and as a fixed-intersection mean over the 123 cases passed by all settings.The fixed intersection prevents differing coverage from changing the performance denominator.
- Knowledge settings: Cold start, Ascend memory, K3 memory, and combined memory isolate source-stack reuse, target-side learning, and their combination.All settings use GLM-5.2 and identical case definitions.
Retargeting Results
Retargeting results show that source and target experience broaden accepted solutions on K3, but target-side experience is more reliable for performance and the combined setting remains below aggregate PyTorch parity.
- Baseline: 1.004× is the measured FlagGems baseline geometric-mean speedup, with parity reached in 76 of 204 cases.Performance varies widely across operators, from large matmul gains to below-parity elementwise and reduction kernels.
- Retargeting Results: 123 cases is cold-start coverage at 60.3%, compared with 170 cases and 83.3% using Ascend-derived memory.The latter adds 47 cases and 23.0 percentage points, with all cold-start passes retained.
- Retargeting Results: 0.145× versus 0.388× is Ascend memory’s fixed-intersection geometric mean relative to cold start, a 62.5% decrease.Source experience transfers feasibility information more reliably than performance choices.
- Retargeting Results: 0.514× is K3 memory’s fixed-intersection geometric mean, while combined memory reaches 0.786× and 188 passes, improving the mean by 1.530× over K3 memory.The combined setting is 2.027× cold start on the same 123 cases.
- Operator coverage: 54.2% to 75.0% is matmul family coverage with Ascend memory, alongside gains from 63.3% to 81.7% for activations, 61.4% to 84.3% for elementwise kernels, and 58.0% to 88.0% for reductions.The improvements span multiple kernel families rather than concentrating only in matmul.
- Operator coverage: K3 memory reaches full coverage for bmm, exp, sub, mul, div_tensor, sum, and argmax; combined memory additionally completes abs, neg, and rsqrt.The combined aggregate remains below parity, indicating adaptation rather than a solved backend.
Implications for KOPE’s Generality
KOPE’s optimization and memory loop is portable across toolchains, but cross-hardware reuse does not guarantee performance transfer. Target-side feedback is needed to refine portable experience for a new backend.
- Interface portability: KOPE executes its optimization and memory loop in a second, heterogeneous toolchain without requiring an AscendC-specific memory node type.The interface consumes candidate kernels and receives compilation, acceptance, and latency feedback while recording decisions, outcomes, and provenance.
- Cross-hardware reuse: 23.0-point coverage gain shows prior experience remains operationally useful after a hardware and programming-model change.This evidence concerns feasibility reuse rather than a claimed speedup.
- Cross-hardware reuse: Fixed-case performance loss shows feasibility knowledge and performance knowledge do not transfer equally across hardware.The results therefore separate coverage reuse from performance transfer.
- Target-side learning: Target-side feedback revises a portable initial prior, supporting KOPE as a retargetable, general-purpose workflow rather than a hardware-independent speedup method.K3 memory and combined memory improve coverage and fixed-case performance over cold start.
Conclusion
The measured RISC-V study supports KOPE’s retargetability across a different hardware and software stack while narrowing the generality claim. It shows useful cross-hardware reuse and target-side performance gains, not hardware-independent speedup.
- Conclusion: KOPE’s RISC-V study spans different hardware, kernel language, compiler path, and evaluator, strengthening its retargetability claim.The study executes KOPE’s optimization and memory loop on a second stack.
- Conclusion: Ascend memory raises pass coverage from 60.3% to 83.3% but reduces fixed-case performance.This separates useful cross-hardware feasibility reuse from performance transfer.
- Conclusion: K3 memory and combined memory raise both coverage and performance, reaching 92.2% coverage and a 0.786× fixed-intersection geometric mean.These results provide the strongest performance evidence in the measured RISC-V study.
- Conclusion: Figure 4 organizes passed/total cases by operator and knowledge setting, with color encoding pass rate across four operator families.The archived pass sets are nested across the settings in these runs.