PaperlayerView exact revision on arXiv

Source-linked AI summary

CANNBot-CAKE: Collaborative Agentic Kernel Evolution for Ascend NPU

Deng Pan, Xue-Jian Gao, Junyi Fan, Bin Du, Jiasheng Li, Xinxian Chen, Nok Yan Chang, Xinlin Li, Lingchao Zheng, Jun Li, Yuwei Fan

arXiv:2610.31034v1cs.AI

TL;DR

AscendC kernel generation remains difficult amid scarce public programming data and limited model success, despite progress in GPU kernel generation. CAKE combines deployable-operator workflows, adaptive routing, and reviewed knowledge reuse; across 53 Ascend 950PR operators it achieves full standard-set coverage and 98.28% hidden-set coverage.

  • Problem

    Public AscendC data are sparse, and general-purpose models have reported very low success on complex AscendC kernels, limiting automated kernel development.

    Source: §1 · p. 1

  • Method

    CAKE generates and evolves deployable Ascend operators through adaptive multi-DSL routing, separate Generation and Evolution modes, and measured-run knowledge reuse.

    Source: §1 · p. 2

  • Results

    Across 53 operators on Ascend 950PR, CAKE achieved full standard-set case coverage and 98.28% hidden-set case coverage.

    Source: §8 · p. 17

  • Takeaways & Limitations

    CAKE organizes operator engineering around deployable artifacts, uses accumulated evidence to select tracks, and carries lessons across runs under human review.

    Source: §8 · p. 17

    EngramGateFusion local and official scores are not directly comparable because the devices differ and round acceptance used a local oracle.

    Source: §4.4 · p. 12

Abstract

from arXiv · show

Efficient custom operators are essential for modern deep learning workloads, yet developing high-performance kernels for Ascend NPUs remains slow, expert-intensive, and difficult to automate. Although AI agents now approach expert performance on GPU kernel benchmarks, current models remain less capable of generating AscendC kernels for NPUs, potentially reflecting the relative scarcity of relevant training data. We thus present CAKE, a CANN-native system for deployable custom operators that combines adaptive multi-DSL routing, separate Generation and Evolution modes, and reviewed knowledge reuse through CAKE Book and a Self-Evolving Agent. On 53 CANN-Bench v1.1 operators on Ascend 950PR, CAKE's aggregated results achieve $100\%$ case pass rate on the standard set (1,060/1,060) and $98.28\%$ on the hidden set (4,167/4,240), with an overall mean score of 78.63 out of 100 and an average performance speedup of $1.334\times$ over the baseline. Evolution then pushes correct operators further: EngramGateFusion cuts aggregate device time by $46.6\%$ over 20 rounds and roughly halves the mean roofline gap. These results offer a practical path toward closing Ascend's AI-for-kernel capability gap.

1 Motivation and System Overview

CAKE addresses the scarcity of AscendC kernel expertise with a workflow for generating, validating, and evolving deployable CANN operators. Its routing, measured-run feedback, and reviewed knowledge reuse distinguish it from systems that automate only one slice of operator engineering.

1 Motivation and System Overview

  • General-purpose models have had very low success on complex AscendC kernels, while hand-written kernels demand scarce expertise and kernel quality sets end-to-end performance.Published evaluations report near-zero zero-shot success on complex AscendC kernels and one-shot pass rates at or below 2.1%, compared with 44.2–52.6% for CUDA.Source: §1 · p. 1
  • CAKE turns a natural-language request or AscendC baseline into an inspectable, deployable operator through separate Generation and Evolution modes.Its pipeline retains artifacts from the operator contract and reference through DSL, host tiling, device code, package, Python binding, and evaluation record.Source: §1 · p. 2, Figure 1 · p. 3
  • The router gates infeasible tracks and scores four code-generation tracks using operator structure and prior evaluations, while the Skills harness supplies reusable lifecycle tools.A self-evolving agent mines measured evaluation runs to update skills, routing, repair templates, and CAKE Book.Source: §1 · p. 2, §1.1 · pp. 4–5

1.1 Related Systems

  • Prior Ascend systems address slices such as model training, DSL transcompilation, kernel optimization, or curated search knowledge, whereas CAKE targets the deployable-operator chain.AgenticCANN is the closest comparison, but its knowledge is authored by experts ahead of time rather than accumulated from measured runs.Source: Table 1 · p. 4, §1.1 · p. 4
  • CAKE Book is an interlinked, cross-user wiki whose measured-run entries become canonical only after human review and on-device verification.This differs from AgenticCANN’s expert-authored taxonomy; CAKE targets reusable lessons accumulated from measured runs.Source: §1.1 · p. 4

2 From Intent to Deployable Operator

CAKE treats operator development as a staged lifecycle whose artifacts and validation records make failures diagnosable and successful operators deployable. Its typed Skills harness binds tools to lifecycle stages, checks outputs, and records execution and recovery.

2 From Intent to Deployable Operator

  • CAKE organizes operator development into stages with explicit inputs, outputs, and validation checks, producing artifacts that support failure diagnosis before the workflow proceeds.A run can stop at a failed intermediate stage while preserving evidence, or leave a package, binding, test record, profile, and provenance trail.Source: §2 · p. 5
  • The operator contract records identity, semantics, I/O constraints, attributes, and cases; a PyTorch reference and stateless functional form support correctness checks and routing.The contract is the source of truth for downstream signatures, while the reference defines expected outputs.Source: §2 · p. 5
  • The router filters tracks by operator and device feasibility, then the selected track lowers the functional form into host tiling and device code for packaging, binding, and evaluation.Evaluation checks precision and on-device timing; failures route back to the relevant stage, while correct but slow candidates enter profiling-guided repair or Evolution mode.Source: §2 · p. 5

2.1 The Skills harness

  • A skill is a typed lifecycle tool with declared input, output, and admissibility conditions; the orchestrating agent currently declares more than thirty skills.New capabilities can be added as skills without editing the agent itself.Source: §2.1 · pp. 5–6
  • Skills bind to lifecycle stages, and machine-checked gates prevent invalid artifacts from advancing while traces and bounded recovery record and address failures.Repeated failure modes must be pushed back into lowering or the project skeleton, and a stage that fails its gate produces no commit.Source: §2.1 · p. 6

3 Adaptive Routing and Generation

CAKE routes operators among feasible code-generation tracks using operator features and evaluation history. Its routing loop records decisions and outcomes so later choices can reflect updated track estimates and skills.

Routing architecture

  • The router combines a feasibility gate, feature extractor, scoring layer, and append-only observation ledger to estimate track suitability from prior evaluations.Track performance can vary with operator structure and current code-generation skills.Source: §3 · p. 6

Feasibility gate

  • The feasibility gate removes unsupported track–chip–operator combinations, allows forward 910B-to-910C compatibility, and abstains when no feasible track remains or the device is unrecognized.On 950 vector and cube–vector operators, only DSL-C is feasible; SIMT operators are accepted only by the SIMT track.Source: §3 · p. 6

Operator features and difficulty

  • Operator features encode eight dimensions, while calibrated difficulty is an operator-level property that cancels when tracks are compared in the scalar regime.Known operators use vectors derived from their reference AST; new operators normally receive heuristic vectors.Source: §3 · p. 7

Track ability and routing score

  • DSL-A’s raw mean achieved score is 0.576 versus DSL-C’s 0.471, but their ability estimates are statistically indistinguishable: −0.95 ± 0.38 versus −0.99 ± 0.36.The raw 0.105 difference alone does not establish a track-ability advantage under the model.Source: §3 · p. 7
  • CAKE models normalized evaluation scores with per-track ability estimates, then fits dimension-specific abilities because a scalar ability would rank tracks identically for every operator.When observations are insufficient for the free parameters, the router falls back to the scalar form; fitted profiles differ most sharply on data movement and blocking.Source: §3 · p. 7

A worked decision

  • On 910B, data-movement features flip the preferred track between gelu and gather, while four worked decisions select three different tracks.The table reports predicted normalized scores by track and identifies the dimension deciding each ordering.Source: §3 · p. 7, Table 2 · p. 8

Evaluation ledger and skill versions

  • The append-only ledger stores evaluation scores and skill fingerprints; superseded skill-version observations count at half weight, while new evaluations receive full weight.Infrastructure failures, reward-hacking candidates, and internally inconsistent submissions remain recorded but are excluded from ability estimation.Source: §3 · p. 8

Exploration and routing loop

  • The exploration heuristic makes alternatives eligible using p(1−p), selects among them by the highest optimistic upper bound, and can bonus tracks with fewer current-version evaluations after skill updates.Each completed evaluation updates subsequent routing; saved candidate snapshots and timeline entries allow past decisions to be reconstructed.Source: §3 · p. 8

4 Evolutionary Optimization

CAKE-evo uses profiling to guide iterative optimization and turns validated improvements into reusable rules. In the EngramGateFusion case study, aggregate device time fell substantially, although local and official scores are not directly comparable.

4 Evolutionary Optimization

  • Generation mode provides a working implementation, while Evolution mode explores tiling, buffering, and fusion choices through parallel candidates and on-device evaluation.Profiling identifies bottlenecks for each round’s optimization targets.Source: §4 · p. 8

4.1 Profiling-driven analysis layer

  • A roofline model estimates an optimistic Tmin from shape and device limits, then uses the dominant pipeline to select bottleneck-driven variants and fusion recommendations.It treats reads and writes as overlapping channels and models output writes as L2-fill-first before HBM traffic; headroom near one indicates little modeled improvement remains.Source: §4.1 · p. 9

4.2 The evolutionary loop

  • Each round generates variants for the identified bottleneck, evaluates and tiers them by speedup, and seeds the next round from those tiers; AIC guides candidates and PMR mines validated rules.Search stops when runtime approaches the modeled floor or the round budget is exhausted.Source: §4.2 · p. 10, Figure 5 · p. 11

4.3 Strategy knowledge

  • The strategy library spans kernel-level dtype, performance, and accuracy changes, plus fusion and host-level shape-driven dispatch; profiling maps bottlenecks to candidate strategies.PMR distills validated performance improvements into reusable rules.Source: §4.3 · p. 11, Table 3 · p. 11

4.4 Case study: evolving EngramGateFusion without an official kernel baseline

  • Across 20 cases, aggregate device time fell from 1451.7 µs in round 1 to 775.7 µs in round 19, a 46.6% reduction or 1.87× within-system speedup.The packaged kernel passed all 20 official grading cases; the campaign reached its 20-round budget, not the roofline stopping threshold.Source: §4.4 · p. 12
  • Local and official scores are not directly comparable because the development device has 48 vector cores while the grading device has 40, changing shape-dependent dispatch.The first official submission failed 12 cases; the final 20/20 result refers to the package repaired in round 20.Source: §4.4 · p. 12
  • EngramGateFusion revisions targeted fixed vector-call overhead, reduced its convolution ring buffer from K times dilation to K rows, and later removed a per-row multiplication.The campaign used 36 generated variants and 1,037 LLM calls across 20 rounds; 186 calls were for final verification and packaging.Source: §4.4 · pp. 12–13
  • Figure 6’s fusion indicator rose from 45.5× to 98.2× as the mean kernel-to-roofline ratio fell from 25.7× to 12.6×; the line is only a trajectory indicator.The indicator mixes eager wall time with device kernel time; a separate same-basis measurement reports a final mean end-to-end speedup of 29.5× (geometric mean 27.6×).Source: Figure 6 · p. 13

5 CAKE Book and System-Level Self-Evolution

CAKE pairs operator-level Evolution with a system-level loop that captures run evidence, reuses evolving knowledge, and promotes stable lessons through human review.

5 CAKE Book and System-Level Self-Evolution

  • CAKE Book and the Self-Evolving Agent turn run experience into reviewed system improvements, making CAKE less dependent on a single model’s prior knowledge.The loop gathers compilation failures, profiling signals, successful optimizations, and human corrections as reusable system memory.Source: §5 · p. 13

5.1 CAKE Book as Agent-Native Reference

  • CAKE Book combines dynamic, version-aware experience with retrieval and updates, but the evaluation does not isolate its causal effect on success or performance.Retrieved context is ranked by relevance, evidence, compatibility, freshness, and review status; updates retain provenance and await review before becoming trusted references.Source: §5.1 · pp. 13–14

5.2 Run Evidence from the CAKE CI Platform

  • The CI Platform exposes run artifacts and groups runs by operator and execution attributes, supporting reusable evidence and regression detection.Records include code, build and evaluation logs, profiles, traces, and failed attempts as well as successful paths.Source: §5.2 · p. 14

5.3 Self-Evolving Agent

  • The Self-Evolving Agent monitors runs, records knowledge updates, and proposes repository changes, while maintainers retain merge authority.CAKE Book is the fast memory; reviewed static references, router rules, templates, and Advisor rules are stabilized memory.Source: §5.3 · p. 14

5.4 Reference Promotion and Governance

  • Promotion requires repeated usefulness, concrete run evidence, clear hardware and software scope, and human review before stable lessons enter shared references.The agent links relevant runs and proposes updates; operator-level optimization supplies observations, while repository review determines which lessons enter the shared workflow.Source: §5.4 · p. 15

6 Evidence from CANN-Bench and Operator Campaigns

On Ascend 950PR, CAKE achieves high standard and hidden case coverage across 53 operators; router records show operator-aware track selection, while CAKE Book records document cross-task knowledge reuse.

6 Evidence from CANN-Bench and Operator Campaigns

  • CAKE passes 1,060/1,060 standard cases (100%) and 4,167/4,240 hidden cases (98.28%), with an overall mean score of 78.63 out of 100.All 53 operators compile in both case sets; all pass every standard case, while 38 pass every hidden case.Source: §6 · p. 15
  • Standard-set case pass rate is 100% at every difficulty level; hidden-set pass rates range from 100% to 94.84%, while mean scores fall to 74.11 and 63.34 at level 4, respectively.Figure 7 groups the 53 operators by difficulty level and shows mean operator score alongside case pass rate.Source: §6 · p. 15, Figure 7 · p. 16

6.1 Router Observations

  • DSL-A’s raw mean exceeds DSL-C’s by 0.105, yet adjusted abilities are statistically indistinguishable; four worked decisions select three tracks as operator features change rankings.The records show choices depend on data movement, blocking, and scheduling/reduction features rather than raw mean alone.Source: §6.1 · pp. 15–16

6.2 Operational Evidence for CAKE Book

  • CAKE Book records show campaign diagnostics retained with provenance and review state, then made available as task-time context for later tasks.These operational records complement the end-to-end results by documenting knowledge capture, review, and subsequent availability.Source: §6.2 · p. 16

7 Scope and Future Directions

The paper distinguishes the 950PR aggregate benchmark from Router and EngramGateFusion evidence on Ascend 910B/910C, and identifies extensions and matched evaluations as future work.

7 Scope and Future Directions

  • The aggregate CANN-Bench results use Ascend 950PR, whereas Router examples and the EngramGateFusion campaign retain Ascend 910B/910C settings.Appendix A documents Ascend 950 programming paths and generation logic.Source: §7 · p. 16
  • Future work will extend Evolution to more operator families and evaluate routing and retrieval configurations under matched task sets and run budgets.The planned evaluation also follows knowledge from CI diagnostics through reuse and reviewed reference updates.Source: §7 · p. 16

8 Conclusion

CAKE treats deployable custom operators as the unit of Ascend engineering and combines adaptive routing, on-device Evolution, and reviewed cross-run knowledge reuse.

8 Conclusion

  • CAKE organizes operator engineering around deployable custom operators with contracts, references, lowering artifacts, packages, bindings, evaluation records, and profiles.Adaptive routing draws on accumulated operator evidence; Evolution searches within one operator contract using on-device feedback; CAKE Book and the Self-Evolving Agent carry lessons across runs under human review.Source: §8 · p. 17
  • Across 53 operators on Ascend 950PR, CAKE achieves full standard-set case coverage and 98.28% hidden-set coverage; EngramGateFusion demonstrates Evolution improving an already-correct operator.The EngramGateFusion improvement uses on-device feedback.Source: §8 · p. 17

A Ascend 950 Vector: Regbase SIMD and SIMT

The Ascend 950 shares its Cube primitives with the 910B, but its vector side adds Regbase and SIMT paths. Regbase targets regular, latency-sensitive kernels; SIMT targets data-dependent memory access or control flow.

A Ascend 950 Vector: Regbase SIMD and SIMT

  • Table 5 contrasts the operand models of membase and Regbase across data residency, addressing, masking, and transfer operations.Source: Table 5 · p. 19
  • The 950 adds two vector-side paths: register-level SIMD Regbase for regular, latency-sensitive kernels and SIMT for data-dependent access or control flow.The 950 shares its Cube primitive set with the earlier 910B.Source: A Ascend 950 Vector: Regbase SIMD and… · pp. 18–19

A.1 Register-level operand model

The 950 Regbase path programs a register-level MicroAPI through @vf micro-functions, keeping vector data in registers rather than UB-resident LocalTensors. Its register operations use fixed-size batches, explicit masks, and raw-byte addressing.

A.1 Register-level operand model

  • RegTensor holds a fixed VREG_SIZE-byte batch, and a MaskReg predicate carries active lanes to each operation instead of a count argument.Source: §A.1 · p. 19
  • Regbase gather and scatter use raw byte offsets without bounds checks, while LoadUnAlign and StoreUnAlign support unaligned UB-to-register transfers.Source: §A.1 · p. 19
  • Operations without a supported register-level lowering, including widening Mull or reading back a Squeeze count, remain on membase.Source: §A.1 · p. 19
  • Regbase runs @vf micro-functions on RegTensor operands held in vector registers, avoiding the UB round trips of membase operations.This shifts the throughput limit from UB bandwidth toward instruction issue, widened by dual issue.Source: §A.1 · p. 19

A.2 SIMT block resources

Each Ascend 950 AI Vector core exposes a SIMT block with up to 2,048 threads in 32-thread warps, supporting warp-level collectives, atomics, and block-wide synchronization.

A.2 SIMT block resources

  • Each AI Vector core supports up to 2,048 SIMT threads in 32-thread warps, with warp shuffle, vote, reduction, and 11-operation atomics.Source: §A.2 · p. 19
  • SIMT threads access a DataCache partition of Unified Buffer capped at 128 KB, with the remaining UB reserved for SIMD.Source: §A.2 · p. 19
  • SIMT blocks synchronize through block-wide barriers.Source: §A.2 · p. 19

A.3 Path selection

CAKE routes operators among SIMT, Regbase, and membase according to access patterns, DSL coverage, and supported operations. For Regbase and SIMT, it generates kernels and uses simulator signals or hardware-resource-aware programming models to guide validation and implementation.

A.3 Path selection

  • On the 950, data-dependent non-contiguous access selects SIMT; otherwise, DSL coverage and the known gap list determine Regbase or membase fallback.The SIMT path further requires irregular, data-dependent access that is non-contiguous; supported contiguous scans can remain on SIMD.Source: §A.3 · p. 19, Figure 8 · p. 20, §A.5 · p. 21

A.4 Regbase code generation

  • Regbase is attempted for register-only micro-instructions, unaligned UB access, or small, high-iteration kernels, only when DSL-C supports all required operations.Membase remains the default.Source: §A.4 · p. 20
  • Regbase emits an @vf kernel with register and mask allocation, CeilDiv(len, VREG_LANES) iterations, and a mask for the final partial iteration.Templates cover lookup-table/gather, compaction/filtering, and histogram/scatter; migration proceeds region by region with bit-exact checks, leaving unsupported regions in membase.Source: §A.4 · p. 20
  • CAKE checks Regbase conversions for fewer same-pipe barriers, a lower load/store ratio, and fewer cannsim wall cycles than membase on the same case.Cross-pipe barriers remain, and per-pipe splits identify the dominant pipe.Source: §A.4 · p. 21
  • Without 950 hardware, cannsim or msprof supports preliminary simulator-only validation.Source: §A.4 · p. 21

A.5 SIMT code generation

  • Pure SIMT supports the full SIMT intrinsic set and suits fully irregular operators, but cannot call the AscendC Micro API.It uses a __global__ kernel and can write output directly to global memory.Source: §A.5 · p. 22
  • The SIMT VF bridge dispatches per-thread compute from an __aicore__ entry and is the default for operators fitting the standard CANN host/device contract.Examples include gather, cumulative sum, and top-k with warp reductions.Source: §A.5 · p. 22
  • Hybrid SIMT+SIMD stages irregular reads in DataCache before vectorizable arithmetic, but its benefit depends on staging and synchronization costs.For fused reductions, high-level AscendC APIs in the __aicore__ entry handle reductions that __simd_vf__ cannot compile.Source: §A.5 · pp. 22–23
Loading 2610.31034v1…