AscendKernelGen：基于大语言模型的NPU算子内核生成系统性研究

arXiv PDF 直链​
https://arxiv.org/pdf/2601.07160v2.pdf ## 系统设计、数据集构建、微调与全面评测

曹新子¹³*、翟建阳¹³*、李鹏飞²*、胡志恒²*、严岑²、穆炳旭²、方广欢²、佘斌²、李佳宇²、苏一涵²、陶东阳²、黄先松¹、徐帆¹、杨飞雕¹、陆瑶¹、王昌东³、卢宇彤³、薛伟成¹†、周斌¹†、田永鸿¹⁴†

  1. 鹏城实验室
  2. 华为
  3. 中山大学
  4. 北京大学
    通讯作者：xuewch@pcl.ac.cn，senosy@gmail.com，yhtian@pku.edu.cn
    arXiv预印本编号：2601.07160v2 [cs.AI]，2026年4月17日

摘要

为满足持续增长的算力效率需求，神经网络处理器（NPU）已成为现代AI基础设施的核心组件。但想要充分释放NPU硬件性能，必须基于厂商专属领域专用语言（DSL）开发高性能计算内核，该工作对硬件专业知识要求极高，人力成本巨大。
尽管大语言模型（LLM）在通用代码生成领域表现亮眼，但在NPU场景下，受限于严苛的硬件约束与领域训练数据稀缺两大问题，生成效果极差。本文前期验证显示，现有主流通用大模型几乎无法生成可正常运行的昇腾NPU复杂内核，任务成功率趋近于0。

针对上述痛点，本文提出AscendKernelGen——一套面向NPU内核开发的生成-评估一体化框架。本文主要创新包含三部分：

  1. 构建Ascend-CoT高质量数据集：收录真实算子工程实现，附带完整思维链推理标注；
  2. 提出KernelGen-LM领域自适应模型：结合有监督微调与带执行反馈的强化学习完成训练；
  3. 设计NPUKernelBench综合评测基准：从编译通过性、功能正确性、硬件性能三个维度，分层评测不同复杂度内核代码。

实验结果表明，本文方案大幅缩小了通用大模型与硬件专用代码生成之间的能力鸿沟：对于二级复杂内核，采样10次时编译通过率从0%提升至95.5%，功能正确率从完全失效提升至64.3%。该结果证明，领域专属推理数据与严谨的硬件执行评估机制，是自动化生成加速器专用代码的核心关键。
本研究开源项目AscendKernelGen与评测基准NPUKernelBench均已公开。

1 引言

深度学习技术飞速发展，彻底重构了现代计算架构。为匹配持续上涨的算力与吞吐需求，领域专用加速器成为主流解决方案，其中华为昇腾系列NPU凭借深度神经网络负载高性能优势，在AI基础设施中占据核心地位。
但加速器的实际算力上限，不只由硬件本身决定，更依赖底层计算内核的实现质量。因此，开发适配硬件、高性能的内核程序，是充分挖掘NPU算力的必要前提。

然而人工编写内核难度极高：昇腾NPU内核开发依赖昇腾C（AscendC）这类厂商专属DSL，要求开发者精通硬件架构细节，包括分层存储精细化调度、分块计算策略设计、计算与数据搬运重叠的异步流水线编程、显式调用向量（SIMD）与矩阵（Cube）运算单元等。整套技术栈学习门槛极高，人工开发耗时久、成本高、极易出错，成为AI应用在NPU上快速迭代部署的核心瓶颈。

从更宏观的视角看，该痛点暴露了现有代码生成方案的底层缺陷：通用编程知识不足以支撑高度专用、绑定硬件的代码生成任务。
学界与工业界借鉴大模型通用代码生成的成功经验，尝试用LLM自动化内核开发，但该路线存在核心障碍：硬件专用DSL所需知识与通用编程语言完全割裂，包含严格API约束、架构绑定语义、性能导向优化范式；同时该领域高质量训练数据极度稀缺，导致通用大模型完全不具备处理这类任务的能力。

本文前期评测（表1）从实证层面验证了该缺陷：零样本场景下，即便是通义千问3等顶尖通用大模型，生成昇腾C内核的效果极差。模型频繁虚构不存在的API（例如AscendC::Softmax），或是严重误用底层硬件接口，编译失败率极高；针对二级、三级复杂内核，执行成功率直接降至0%。这说明，未经领域适配的通用大模型，完全无法处理有实际业务价值的复杂NPU内核开发任务。

表1 通用大模型在NPUKernelBench基准上的零样本生成效果

模型 任务难度 编译通过率 编译后功能正确率
通义千问3-8B 一级（简单算子） 8.22% 1.08%
二/三级（复杂算子） 1.39% 0.00%
通义千问2.5-Coder-7B 一级（简单算子） 9.19% 0.47%
二/三级（复杂算子） 0.40% 0.00%
Llama3.1-8B 一级（简单算子） 23.97% 0.69%
二/三级（复杂算子） 19.44% 0.00%
Mistral-7B 一级（简单算子） 0.00% 0.00%
二/三级（复杂算子） 0.00% 0.00%

上述结论说明，想要让大模型生成高质量NPU内核，必须完成两大核心改进：

  1. 让模型深度掌握硬件专属编程范式，吃透API约束与硬件架构特性；
  2. 搭建专业可靠的评估框架，不仅校验代码能否编译，还需验证功能正确性与运行性能——三者均是内核代码的核心评价指标。

为此，本文提出AscendKernelGen，面向NPU内核开发的生成-评估一体化框架。本文三大核心贡献如下：

  1. 面向内核生成的领域专属推理数据集：构建并开源Ascend-CoT高质量数据集，提炼底层NPU内核编程所需完整结构化推理流程，数据源自真实工程内核实现，附带流水线搭建、同步逻辑、数值计算推理等详细标注；
  2. 底层代码生成领域自适应后训练方案：基于基座大模型，提出面向NPU内核合成推理难点的领域自适应后训练策略，得到KernelGen-LM模型；相比通用代码大模型，编译通过率、功能正确率均大幅提升；
  3. NPU内核综合评测基准：设计NPUKernelBench评估框架，从编译、功能、性能三个维度系统化评测生成内核；双路径评估机制可分别静态尺寸优化、动态尺寸鲁棒性开展精细分析，为后续加速器专用代码生成研究提供标准化评测工具。

实验结果证明本文方案有效性：在Ascend-CoT数据集完成领域自适应微调后，模型内核生成质量全面提升，尤其是零样本场景下完全无法解决的复杂算子；各级难度任务的编译通过率、功能正确率同步上涨，证明领域专属推理监督，是填平通用大模型与硬件内核编程鸿沟的关键。

本文组织结构如下：第2章梳理大模型代码生成、传统硬件内核优化、大模型内核生成、代码生成评测基准四大方向相关工作；第3章介绍NPU内核开发的编程抽象与硬件约束；第4章整体介绍AscendKernelGen框架，包含内核生成流水线、硬件验证评估链路；第5章详述训练数据集构建，分为文档思维链、代码思维链、通用思维链三类数据；第6章讲解面向NPU内核的大模型微调方案；第7章介绍评测基准NPUKernelBench，基于真实昇腾硬件完成编译、正确性、性能全流程验证；第8章介绍实验配置，深度分析实验指标；第9章总结全文并展望未来研究方向。

2 相关工作

本文研究横跨两大领域：大模型代码生成、硬件专用内核优化。本章从四个维度梳理现有研究：通用大模型代码生成、传统硬件内核优化、基于大模型的内核生成、代码生成评测基准。

2.1 通用大模型代码生成

大语言模型在Python、C++、Java等主流编程语言的通用代码生成任务上取得巨大突破，代表性工作包括OpenAI Codex、AlphaCode、Code Llama、通义千问2.5/3代码系列、CodeRL等，可实现代码补全、重构、算法竞赛解题等功能。

尽管成果显著，但通用大模型在包含专用API、强语义约束的领域代码生成任务中存在明显短板。硬件专用编程语言的开发要求与通用软件开发存在本质差异：必须手动管理分层存储、架构绑定执行语义、性能优先的正确性约束。面对这类场景，大模型极易生成无法编译、语义错误的代码，因此亟需领域专属数据与训练策略。

2.2 传统硬件内核优化

在大模型普及前，专用硬件内核优化主要依靠编译器、自动调优两条路线。

  1. 自动调优框架：TVM、Ansor等工具在调度与优化空间内搜索最优方案，适配多类硬件后端；
  2. 高层抽象DSL：Halide、Triton、MLIR等将算法描述与执行调度解耦，大幅降低GPU等加速器内核开发门槛；TensorFlow XLA、PyTorch 2.0基于编译器流水线，将高层计算图转为后端专用优化内核；
  3. 多面体编译：Tensor Comprehensions、Tiramisu、AKG（昇腾NPU专用）系统性建模并优化循环变换；
  4. 算子搜索与代数变换：TASO、PET、EINNET等；
  5. 平台专用高性能库：CUTLASS、Liger-Kernel提供手工优化内核，适配特定硬件与负载。

上述方案虽能实现高性能，但普遍存在缺陷：依赖大量专家知识、搜索/开发成本极高、跨架构泛化能力弱；通用编译器难以适配闭源、接口受限的硬件平台，平台专属库可移植性差。因此学界开始探索更灵活、基于学习的内核生成方案。

2.3 基于大模型的硬件内核生成

受大模型通用编程能力启发，近年研究开始探索用LLM自动生成硬件内核，现有工作主要集中在CUDA GPU、TPU等成熟生态。
早期方案依靠提示词、上下文学习直接生成内核，代表工作CUDA-LLM、AI CUDA Engineer；后续引入有监督微调做领域适配，KernelLLM证明小型领域专用模型即可生成具备竞争力的GPU内核。

近期研究结合强化学习与执行反馈迭代优化内核：Kevin、CUDA-L1、AutoTriton、TritonRL等方案证明，有监督初始化+执行导向优化可显著提升代码正确率与性能。同时，智能体、多智能体框架被用于自动化内核优化，完成规划、调试、空间搜索全流程，例如EvoEngineer、STARK、CudaForge；也有研究基于检索、图结构推理优化内核。

但现有基于LLM的内核生成工作几乎全部面向CUDA、TPU这类文档完善、工具链成熟、训练数据充足的平台；针对昇腾这类采用闭源编程模型的新兴NPU，基于大模型的内核生成研究仍存在大量空白。

2.4 代码生成评测基准

通用代码生成领域已有标准化评测集，例如HumanEval、MBPP，但这类基准仅关注功能正确性，不适用于硬件内核生成——内核任务中编译通过性、硬件约束合规性、运行性能同等重要。

为填补该空白，KernelBench、MultiKernelBench等硬件导向基准新增编译、性能评估维度，但普遍采用单链路评测范式，要求生成完整主机-设备耦合代码，对简单算子冗余度高，也无法灵活评估不同抽象层级的内核；TritonGym为GPU内核提供交互式工具化评测环境。

整体来看，目前缺少一套适配NPU内核生成、灵活且全面的评估框架，可同时支持独立内核优化、完整主机-设备算子集成评测。搭建该类工具，是可靠衡量专用硬件平台大模型内核生成能力的基础。

3 NPU内核生成的编程抽象与约束

本章梳理底层NPU内核开发的编程抽象，解释该任务为何对大模型生成存在天然挑战。不堆砌硬件细节，重点分析决定代码正确性、性能的程序结构特征。

3.1 结构化内核编程抽象

本文将底层NPU内核定义为一套静态结构化程序，在统一的并行执行模板下，同时描述全局数据划分、异步流水线阶段、显式同步语义。

  1. 并行执行模型：单份内核程序在多个计算单元上复制执行，每个实例根据逻辑块索引处理全局数据的独立分片。代码必须手动计算全局内存偏移、有效数据边界，不能依赖隐式索引；
  2. 异步流水线机制：数据加载、计算、写回分别映射至独立硬件单元，设计目标是时间重叠。但运行时不会自动处理重叠逻辑，内核代码必须通过同步原语显式定义流水线阶段的生产者-消费者依赖，形成静态执行调度；
  3. 底层接口约束：编程接口要求手动指定访存模式、步长、掩码、同步事件，数值计算、数据搬运、控制流高度耦合，编译器不会自动推导优化。

3.2 对大模型代码生成的影响

这套结构化编程抽象带来远超普通代码补全的推理需求，是NPU内核生成难度极高的核心原因：

  1. 长距离语义依赖：内核正确性依赖分块因子、边界尺寸等全局参数，跨越多段流水线，同时控制访存与同步逻辑。生成正确代码需要跨代码段的语义一致性推理，而非局部模式匹配；
  2. 显式同步逻辑推理：异步流水线仅靠手动插入的同步原语完成通信，执行顺序完全由代码定义。同步指令缺失、冗余、顺序错乱都会引发死锁或数据错误，要求模型推理动态执行时序，而非线性指令序列；
  3. 边界敏感数值推理：张量尺寸往往无法对齐硬件最优分块，代码必须通过偏移计算、掩码逻辑处理边界场景。循环边界、索引、有效性判断的微小计算错误，都会导致输出错误或非法访存；
  4. 布局感知表征推理：不同流水线阶段采用适配运算单元的物理数据布局，代码需要在保持张量逻辑含义的前提下处理布局转换。该层物理存储与抽象张量的分离逻辑，无法通过形状变换隐式推导。

4 系统整体框架

本文提出AscendKernelGen（简称AKGen），一套用于研究大模型自动生成底层NPU内核的统一研究框架。本框架并非独立工程系统，而是用于在真实硬件编程约束下，系统化分析内核生成、验证、评估全流程的实验平台。

如图1所示，AKGen由三大紧密耦合模块组成：领域专属推理数据集、内核生成大模型、结构化评测基准。三者形成闭环生成-评估链路，可可控分析模型能力与缺陷。

生成模块

AKGen生成侧由两大组件构成：

  1. Ascend-CoT思维链数据集：解决底层NPU编程数据稀缺问题，收录真实昇腾算子实现，提炼昇腾C内核开发所需结构化推理流程，重点覆盖流水线搭建、同步逻辑、数值计算核心推理范式；
  2. KernelGen-LM模型：基于Ascend-CoT数据集完成领域自适应后训练，专门针对底层NPU内核生成优化。训练目标直接匹配硬件感知代码生成的推理难点：数据分块、异步流水线调度、底层编程接口规范调用。

评估模块

AKGen配套专用评测框架NPUKernelBench，打通大模型概率输出与硬件确定性执行的鸿沟。
NPUKernelBench将评测任务划分为多级难度，从简单逐元素算子到复杂融合算子，可精细分析不同计算模式、编程复杂度下的模型表现。
每一段生成内核都会经过多阶段校验链路：校验编译通过性、对照参考实现验证功能正确性、统计运行时延评估性能。分层设计可隔离各类失效模式，量化超出代码表层合法性的模型能力。本研究当前重点评估正确性与性能指标，框架可灵活扩展静态代码质量分析、深度性能剖析等功能。

整体而言，AKGen提供一套可控实验体系，用于分析大模型如何学习、应用、泛化底层硬件编程知识，昇腾NPU内核作为代表性高难度测试载体。

5 数据集构建

为赋予大模型底层NPU内核生成所需的完整推理能力，本文构建多源融合数据集，融合领域内核知识与通用推理监督。设计目标不只是让模型记住昇腾C内核代码表层写法，而是教会开发者级推理逻辑：分块策略、内存搬运、API约束、硬件执行模型下的正确性保障。

如图1，数据集包含三类互补数据：源自昇腾C官方文档的文档类思维链、从真实昇腾算子提取的内核中心思维链、用于保留泛化能力的通用推理链。原始语料共83916条样本，数据集Ascend-CoT已开源。

原始数据集文本长度呈长尾分布：99.1%样本输入序列长度小于11100token，单样本最大输入长度约111000token；输出侧85.1%样本长度在21900token以内，6.3%介于21900~43800token，最大输出长度219000token。
有监督微调前执行长度过滤预处理：丢弃超过32000token的样本；短样本采用边界隔离打包为定长训练序列，避免拼接样本间互相干扰。预处理后最终SFT训练集包含9955条32000token序列。

5.1 文档类思维链数据

昇腾C开发生态包含大量高度专业化资料，覆盖硬件抽象、编程模型、底层API语义。本文不采用无监督持续预训练，而是采用知识指令范式，将官方文档转化为带显式推理链路的问答对，可控地向大模型注入领域知识。
数据来源分为三类权威文档：《昇腾C算子开发指南》（硬件抽象、编程模型）、《昇腾C API参考手册》（内核/主机/调试接口）、《昇腾C最佳实践》（性能优化策略、常见坑点）。
生成的监督数据不仅覆盖API、硬件概念事实，还包含正确调用、约束满足背后的完整推理逻辑，缓解浅层模式记忆带来的模型幻觉，引导模型系统性推理硬件与编程约束。
数据集同时覆盖硬件抽象、编程语义的概念推理，以及API调用模式、参数约束、典型错误场景的接口推理；每条样本包含自然语言问题、结构化思维链、有依据的标准答案。为保证覆盖度与多样性，严格对齐文档章节，并通过复述增强扩充数据。

5.2 内核中心思维链数据

除文档知识外，正确生成内核需要理解主机侧分块逻辑与设备侧内核执行的交互关系。本文从真实昇腾算子工程实现中提取内核中心思维链数据。

  1. 独立可编译内核文件：生成内核级思维链样本，完整解释内核结构、内存搬运、计算逻辑；每条样本绑定功能描述、完整内核实现、显式推理过程，解释分块假设、缓冲区使用、流水线组织等核心设计决策；
  2. 工业级复合算子：多包含多尺寸、多数据类型、多执行分支，主机分块逻辑与设备内核调用交织，无法直接用于监督。本文将复杂算子拆解为逻辑纯净的主机-内核配对，每组配对对应一套固定分块、尺寸假设的执行场景，生成工程级思维链监督数据；同时推理主机分块参数计算、设备流水线与内存缓冲、分块元数据与内核执行的一致性，显式暴露跨边界依赖关系；
  3. 过滤校验：仅保留自动化验证通过、匹配标准分块配置的思维链样本，保证推理链路逻辑自洽、符合主机-内核执行语义。

5.3 通用思维链数据

仅使用内核领域数据训练会损害模型通用推理能力。为平衡领域特化与通用能力，本文补充高质量开源通用思维链语料，覆盖数学、代码推理、科学问题求解。
统一过滤、归一化处理异构数据源，包含文本标准化、长度过滤、困惑度质量筛选。混合训练后，模型既能适配NPU内核的高度专用推理需求，也保留通用问题求解能力。

6 面向NPU内核的大模型微调方案

为让大模型在严苛编程约束下稳定生成可用NPU内核，本文提出KernelGen-LM领域专用模型，采用有监督微调（SFT）+强化学习（RL） 两阶段优化策略：

  1. SFT阶段：基于精修领域数据集完成NPU适配有监督微调，让模型掌握昇腾C语法、API、内核基础结构；
  2. RL阶段：基于执行正确性反馈的强化学习，进一步优化，让生成内核不仅能编译，还能保证数值计算准确。

6.1 基于错误样本增强的有监督微调

仅依靠静态标准样本微调，模型仍频繁出现编译失败、数值计算错误。本文提出错误样本增强监督策略，将真实执行失败产生的修正数据加入SFT训练，分为API级错误、内核级数值错误两类（图2）。

6.1.1 API级错误修正数据

基于真实昇腾C编译失败日志构建API纠错数据集。输入编译报错日志、对应内核代码、官方参考文档，训练模型定位报错根因并输出修正实现。专门解决大模型生成内核的高频问题：API调用语法看似合规，但违反语义约束、使用规范。完整构建流程见算法1。

6.1.2 内核级数值错误修正数据

大量代码编译通过但数值校验失败，根源是内核深层推理缺陷：内存缓冲逻辑错误、累加顺序不匹配、主机分块元数据与设备执行逻辑不一致。
本文构建真值导向重构数据集（算法2）：筛选编译通过但数值校验失败的内核，配对标准参考实现作为监督目标；生成重构导向思维链，引导模型分析数值错误根源、从零生成修正内核。不同于强化学习，该方案通过样本级硬约束强制满足执行、数值不变性，不依赖标量奖励、偏好对比。
修正样本全部并入SFT训练集，让模型内化数值敏感的执行约束，大幅减少隐性正确性错误；在强化学习前过滤大量无效候选样本，缩小搜索空间，稳定后续策略优化，提升样本效率。

6.2 基于执行偏好的强化学习

经过错误样本增强SFT后，模型已能生成大量可编译、基础数值正确的内核。但同一算子可能存在多份合法内核实现，在内存访问、累加顺序、数值稳定性上存在细微差异，这类细节无法通过确定性静态监督覆盖。
本文引入基于执行偏好信号的强化学习，针对已编译、校验通过的内核做精细化优化，优先强化高质量实现。
流程（图1）：对每个内核任务采样多份可执行候选代码，自动硬件执行校验，基于数值精度、执行正确性构建偏好对；编译+精度校验全通过的内核作为正样本，编译通过但精度失效的作为负样本，形成相对偏好信号，更新模型，持续输出稳定、高精度的内核。
前置错误增强SFT为RL提供支撑：提前过滤大量无效、数值错误候选，大幅缩小策略搜索空间，稳定偏好学习。强化学习仅在合法内核中做精细优化，而非修复致命执行错误，训练收敛更快、效率更高。

7 评测子系统：NPUKernelBench

7.1 基准整体概述

NPUKernelBench是一套端到端评测框架，完整覆盖任务定义、代码生成、编译、硬件执行、量化指标统计全生命周期，用于评估大模型生成正确、高性能NPU内核的能力。
给定标准化内核任务与指令，框架引导大模型同时输出结构化思维链、可执行主机代码与设备内核代码；生成代码在昇腾NPU硬件完成编译运行，对照官方参考实现自动校验功能正确性、统计运行性能。
整套评测完全基于真实硬件行为，强制完整编译与执行流程，可客观、可复现、精细化分析内核生成质量。

7.2 内核任务分层分类设计

评估大模型内核生成能力需要分离多重难度来源，若将所有内核视为同质任务，会掩盖计算结构、接口需求带来的能力差异。NPUKernelBench从两个正交维度划分任务难度：算法复杂度、接口复杂度，形成分层分类评测体系。

算法复杂度（三级分层）
  • 一级任务：简单逐元素、基础算术算子，线性数据流；
  • 二级任务：主流神经网络算子，结构化计算、局部数据复用依赖；
  • 三级任务：存在全局依赖、迭代计算、动态控制流的复杂算子（矩阵乘、TopK、注意力算子等），对并行推理、分层存储调度、控制流代码生成要求极高。
接口复杂度：静态尺寸 / 动态尺寸双路径
  • 静态尺寸：张量维度编译期已知，侧重专用优化、极致吞吐、片上存储利用率；
  • 动态尺寸：运行时张量尺寸可变，要求模型生成鲁棒主机逻辑，完成尺寸推导、分块参数动态计算、运行时边界判断。

7.3 标准化生成接口

为保证评测公平、可复现，NPUKernelBench定义统一提示词规范、输出代码结构规范，打通大模型概率输出与NPU编译执行链路的确定性要求。

7.3.1 提示词组成

提示词由三部分组成：

  1. API描述文件（API_Desc.md）：封装算子功能需求，数学定义、输入输出张量属性、硬件约束；
  2. 代码模板脚手架：主机模板（运行时注册、分块接口）、内核模板（设备侧入口函数），剥离样板代码，强制模型填充合法C++结构，大幅提升代码可编译性；
  3. 角色约束指令：定义模型身份为专业昇腾内核工程师，规定代码规范、输出格式，要求主机、内核代码分块输出，无冗余对话文本。

整套提示词流程固定：先加载API描述，拼接代码模板，前置角色约束指令，形成统一输入上下文，最大化生成合规昇腾C代码的概率。

7.3.2 输出代码规范与双路径评测

框架强制标准化输出代码，支持两条互补评测链路：

  1. 仅设备内核链路：提供预优化主机驱动，仅评测模型生成设备侧计算逻辑的能力，适合一、二级简单算子；
  2. 主机+设备完整链路：要求模型生成完整可部署算子，同时输出主机调度逻辑、设备内核，贴合真实工程场景，覆盖运行时调度、尺寸推导、动态内存管理。

完整链路要求生成三类耦合逻辑：

  • 主机侧代码：CPU运行，负责算子调度、张量尺寸推导、适配NPU的分块参数计算、内核启动维度配置；
  • 设备侧内核：NPU AI核执行，实现并行计算，合规管理片上分层存储、并行同步、双缓冲等流水线重叠优化；
  • 分块结构体：连接主机分块计算与设备内核执行的参数载体。

双路径设计可全面评估两类工业场景：静态尺寸极致优化、动态尺寸鲁棒性。静态路径分析固定尺寸下的分块策略、片上存储利用、计算-访存流水线优化；动态路径检验模型适配运行时多变张量尺寸的能力，分块参数、循环边界动态计算逻辑直接决定代码部署鲁棒性。

7.4 自动化评测流水线

NPUKernelBench不只是测试用例集合，而是高可靠端到端自动化框架，打通代码生成、编译、正确性校验、性能基准四大传统隔离环节。

7.4.1 端到端架构与并发调度

整套评测一键启动，统一配置文件管理，无人工干预，保证多轮实验一致性。
大规模批量评测采用Python多进程高并发执行，并行编译、并行硬件运行，最大化主机CPU与NPU硬件利用率；每条内核评测独立隔离子进程，段错误、非法访存、NPU运行时崩溃等致命错误不会扩散、中断全局任务。
配套稳定机制：超时终止卡死内核/硬件挂起、自动重试处理临时资源波动、运行结束强制回收主机内存与NPU显存，避免资源泄漏，保障超长批量评测稳定运行。

评测遵循三大核心原则：

  1. 正确性优先：必须通过功能校验后，才评估性能；
  2. 性能导向评估：代码合法前提下，衡量NPU硬件资源利用率；
  3. 多级难度分层：从基础实现到高度调优算子逐级提升难度，系统化评估优化能力。

完整自动化流水线分为三阶段，并行调度执行，自动记录全量日志与中间结果：编译评估、正确性评估、性能评估。

7.4.2 编译评估

使用昇腾C编译器处理生成代码，记录编译成功/失败状态，完整日志存档，用于定位编译报错根源。

7.4.3 正确性评估

在多组数据类型、张量尺寸测试用例上执行生成内核，逐元素对比官方真值输出，在预设数值误差阈值内判定是否通过。

  1. 单任务得分：当前任务通过测试用例数 / 总测试用例 × 100；
  2. 层级得分：同一难度下所有任务得分均值；
  3. 综合总分：层级得分加权求和，复杂算子权重更高（一级0.2、二级0.3、三级0.5），鼓励模型优先攻克高难度算子。

静态尺寸任务：所有测试用例对应同一固定输入；动态尺寸任务：单内核必须兼容多套张量尺寸，任意尺寸校验失败都会降低得分，天然衡量代码泛化鲁棒性。

7.4.4 性能评估

仅对功能校验通过的内核评估性能，核心指标为内核运行时延。
每组测试用例预热固定轮次消除冷启动，多次连续运行取平均时延作为性能指标；不只用绝对时延，而是对比厂商官方高度优化内核的参考时延T_ref，计算相对加速比，消除算子规模、张量尺寸、硬件差异带来的干扰，实现跨负载公平对比。

7.5 面向优化的细粒度日志系统

评测日志不只是被动记录工具，更是连接评估与模型迭代优化的反馈链路。全流程结构化存储精准诊断信息：编译器错误码、报错行号、张量数值偏差详情，完整复现每条失效样本。
日志将粗粒度“通过/失败”结果转化为丰富反馈信号，既可人工调优提示词，也可自动化支撑SFT、RL训练：将内核生成拆解为语法合法、编译通过、数值正确、性能达标多个可校验里程碑，提供稠密可解释奖励信号，适配奖励塑形策略，实现训练过程增量信用分配。
NPUKernelBench不只是静态基准集，更是一套评测驱动的训练环境，支撑自迭代内核生成智能体开发。

8 实验与结果

8.1 实验环境配置

模型选型
  1. 数据集构建阶段：DeepSeek-R1生成文档、单文件内核思维链；Gemini 2.5 Pro处理多文件工程级跨依赖思维链；DeepSeek-Reasoner迭代修正思维链样本；
  2. 微调基座：以Qwen3-32B为主干；同时覆盖1.7B/4B/8B/14B/32B、Qwen3-Coder-30B多尺寸模型做缩放实验；微调分为全参数微调、LoRA低秩微调两组对照。
训练配置
  1. SFT阶段：全微调学习率1.25e-6，LoRA微调1.25e-5；余弦衰减，预热比例0.01，全局批次128，权重衰减1e-1，梯度裁剪1.0；
  2. RL阶段：采用DPO直接偏好优化，β=0.1；学习率1e-6，余弦衰减，预热0.1，批次64，共150轮迭代。

8.2 评测基准与指标

全部实验基于NPUKernelBench，共158个代表性内核，覆盖三级计算复杂度，每个算子配套官方参考实现，自动化数值校验。
每条任务调用大模型生成内核代码，在昇腾NPU硬件执行，从三大维度评估：

  1. 编译通过率（Pass@k）：采样k份候选中至少1份编译成功的任务占比；
  2. 执行正确率（Pass@k）：编译通过且数值校验匹配真值的任务占比；
  3. 性能加速比：生成内核时延 / 厂商优化内核参考时延，数值大于1代表性能优于官方基准。

8.3 核心实验结果

对比基座模型、SFT微调、SFT+RL完整链路三阶段生成效果，随训练流程推进，各级难度任务指标稳定提升，复杂算子收益最显著（表7）。

8.3.1 编译与执行正确率

基座Qwen3-32B零样本场景：一级简单算子仅能达到中等编译、执行通过率；二级复杂算子几乎完全失效；三级算子采样数量提升仅小幅改善编译通过率，但执行正确率始终为0，证明编译通过不等于内核逻辑正确。

  • 基座模型平均Pass@1仅7.92%；
  • SFT微调后平均Pass@1提升至26.26%，模型掌握基础API、内存对象、标准计算模板，满足合法内核生成基础条件；
  • SFT+RL强化学习后平均Pass@1进一步提升至33.46%。SFT解决粗粒度语法、编译约束，RL基于执行反馈区分多份合法内核间细微逻辑差异，对复杂算子内存调度、同步、累加顺序类隐性错误改善效果突出。

两阶段训练互补：SFT搭建内核生成基础能力，RL精细化优化执行逻辑，大幅提升复杂算子首次生成正确率。

8.3.2 运行性能加速比

基座模型完全不具备硬件优化能力：一级算子加速比仅0.60倍，二、三级算子无可用合法内核；
SFT微调后模型可生成具备优化能力的内核，二级算子平均加速比1.50倍，超越官方手工优化基准，证明领域思维链数据让模型掌握分块、数据复用、算子融合等硬件优化范式；
叠加RL后二级算子加速比提升至1.86倍，一级算子稳定0.61倍。强化学习带来的收益来源于内存访存顺序、同步指令位置、累加结构的精细化调整，而非大规模代码结构改动。
性能提升在二级算子上最为明显：结构化计算+局部数据依赖，优化空间充足且模型易于推理；一级算子优化空间极小，三级算子全局依赖、复杂控制流仍是难点。整体证明本文SFT+RL流水线不仅能生成语法合法内核，还能产出硬件高效并行逻辑，大量主流算子性能达到甚至超越专家手工实现。

8.4 有监督微调消融实验

8.4.1 模型规模缩放实验

模型参数量提升，一、二级算子编译、执行正确率同步上涨；三级算子难度极高，绝大多数模型执行正确率趋近于0，仅32B规模可实现可观编译通过率，说明超大参数量是处理复杂内核控制流、多层内存依赖的必要条件。

8.4.2 全参数微调 vs LoRA微调对照

基于Qwen3-8B对比两种微调方案：全参数微调全面优于LoRA，平均编译通过率40.29%→55.32%，平均执行正确率13.55%→22.13%；算子复杂度越高差距越大。
内核生成属于高维领域知识学习任务，不只是格式对齐。LoRA低秩分解容量不足，无法完整吸收硬件底层指令、存储约束等复杂语义；全参数更新可完整注入推理逻辑，平均性能加速比0.48倍提升至0.95倍。结论：高精度硬件代码生成任务必须采用全参数微调。

8.4.3 训练数据成分消融

移除内核代码样本后，Pass@1指标断崖式下跌，证明内核代码是核心训练数据；文档、通用推理数据提供辅助增益，对复杂算子提升明显；训练数据规模扩大，生成效果持续稳步上涨，证明数据体量至关重要。

8.5 强化学习消融实验

对比不同负样本构造策略、学习率衰减方案：
以“编译通过但数值校验失败”的内核作为负样本，比“编译直接失败”样本提供更有效的监督，编译、执行正确率全面提升；同等负样本策略下，余弦学习率衰减优于恒定学习率，执行正确率提升显著。
因此本文默认采用编译通过但精度失效样本构建偏好对，搭配余弦衰减学习率。

8.6 错误类型分析

本文统计4000份失败生成样本，拆解错误分布：

  1. API签名/重载错误：51.9%（最高频），函数调用参数顺序、类型、必填项不匹配，混淆同名API语义差异；
  2. 数据类型与转换错误：19.8%，非法硬件数据类型、不兼容类型强制转换；
  3. 变量作用域与生命周期错误：16.4%，变量跨主机/内核上下文未定义、访问失效符号；
  4. 内存与对象使用错误：8.1%，LocalTensor、TPipe等硬件对象调用非法接口、内存地址计算错误；
  5. 语法、宏预处理等小众错误合计3.7%。

结论：模型基本能写出语法通顺代码，绝大多数失效来自底层硬件语义约束，核心瓶颈是精准匹配API规范、硬件专属类型、跨代码段变量可见性推理。

9 总结

本文提出AscendKernelGen一体化框架，用于自动化生成高性能昇腾NPU算子内核。研究证明：大模型处理该任务的核心障碍并非语法陌生，而是缺少硬件约束、分层存储调度、异步流水线同步的结构化推理能力。
本文通过构建Ascend-CoT领域思维链数据集、搭建兼顾编译、数值正确性、硬件性能的NPUKernelBench评测基准，搭配SFT+执行反馈RL两阶段领域自适应训练方案，大幅提升昇腾C内核生成效果，解决通用大模型无法处理复杂算子的痛点。

现有工作仍存在局限，未来三个研究方向：

  1. 优化三级复杂算子生成能力，当前仍是模型主要短板；
  2. 在强化学习奖励中融入性能指标，优先生成低时延、硬件资源利用率贴近硬件上限的内核；
  3. 将整套方法论泛化至其他新兴加速器平台，构建AI驱动软硬件协同设计通用生态。

附录代码说明

1 数值校验标准函数

默认精度校验逻辑

def check_precision(outputs, outputs_new, max_abs_error, max_rel_error):
    # outputs：参考模型输出张量列表
    # outputs_new：大模型生成内核输出张量列表
    outputs_new = [outputs_new] if not isinstance(outputs_new, list) else outputs_new
    outputs = [outputs] if not isinstance(outputs, list) else outputs
    all_abs_diff, all_rel_diff = [], []
    is_accurate = True
    # 逐组张量对比
    for out, out_new in zip(outputs, outputs_new):
        abs_diff = torch.abs(out - out_new)
        rel_diff = abs_diff / (torch.abs(out) + 1e-7) # 加极小值避免除零
        all_abs_diff.append(abs_diff.view(-1))
        all_rel_diff.append(rel_diff.view(-1))
        # 任意元素同时超出绝对、相对误差阈值则判定错误
        if ((abs_diff > max_abs_error) & (rel_diff > max_rel_error)).any():
            is_accurate = False
    all_abs_diff = torch.cat(all_abs_diff)
    all_rel_diff = torch.cat(all_rel_diff)
    return (1 if is_accurate else 0), all_abs_diff, all_rel_diff

自定义精度校验逻辑

def custom_check_precision(param, outputs, outputs_new):
    dtype_str = param.get("dtype", "float16")
    dtype = getattr(torch, dtype_str)
    if dtype == torch.float32:
        # float32高精度阈值
        return check_precision(outputs, outputs_new, max_abs_error=0.00001, max_rel_error=0.00001)
    else:
        # 半精度放宽误差阈值
        return check_precision(outputs, outputs_new, max_abs_error=0.001, max_rel_error=0.001)

2 内核生成提示词完整示例（矩阵乘算子）

API描述文档（API_Desc.md）

aclnnBasicMatmul算子

算子功能

昇腾C内核实现二维矩阵乘法，是深度学习线性层、注意力机制核心组件；输入两个满足矩阵乘法规则的二维张量，输出乘积矩阵。

计算公式

输入矩阵A(m×k)、B(k×n)，输出C(m×n)：
Cij=∑p=1kAipBpjC_{ij} = \sum_{p=1}^{k} A_{ip} B_{pj}Cij​=p=1∑k​Aip​Bpj​
i∈[1,m]，j∈[1,n]，CijC_{ij}Cij​为输出矩阵第i行j列元素。

计算流程与类型转换

为防止大规模累加溢出、保障数值精度，计算采用高精度累加策略：

  1. 输入张量a、b数据类型float16；
  2. 乘累加过程中，中间累加器转换为float32高精度存储；
  3. 全部累加完成后得到float32中间结果；
  4. 最终将float32转回float16作为输出。

接口定义

输入：

  • a：设备侧aclTensor，对应矩阵A，支持float16、二维/多维ND格式；
  • b：设备侧aclTensor，对应矩阵B，支持float16、二维/多维ND格式；
    输出：
  • c：设备侧aclTensor，对应矩阵C，支持float16、二维/多维ND格式；
    属性：无

约束限制

  1. 输入仅支持float16数据类型；
  2. 输入必须为二维矩阵；
  3. 仅支持ND数据排布格式；
  4. A的第二维（列数）必须等于B的第一维（行数）。

主机模板代码

#include "register/op_def_registry.h"
#include "tiling/platform/platform_ascendc.h"
namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext *context)
{
    context->SetBlockDim(platform_ascendc::PlatformAscendCManager::GetInstance()->GetCoreNumAic());
    return ge::GRAPH_SUCCESS;
}
} // namespace optiling

namespace ops {
class BasicMatmul : public OpDef {
public:
    explicit BasicMatmul(const char *name) : OpDef(name)
    {
        this->Input("a")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND});
        this->Input("b")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND});
        this->Output("c")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND});
        this->AICore()
            .SetTiling(optiling::TilingFunc)
            .AddConfig("ascend910_93")
            .AddConfig("ascend910b");
    }
};
OP_ADD(BasicMatmul);
} // namespace ops

内核模板代码

#include <kernel_operator.h>
using namespace AscendC;
extern "C" __global__ __aicore__ void basic_matmul(GM_ADDR a, GM_ADDR b, GM_ADDR c, GM_ADDR workspace, GM_ADDR tiling)

3 各类错误案例对照表

错误分类 典型算子案例 错误说明
API签名/重载错误 Equal、Muls、Arange Muls第三个参数传LocalTensor而非标量；Arange参数数量、类型不匹配；调用不存在的接口
数据类型转换错误 Less、Duplicate float转bool无内置转换指令；uint8_t、int16_t转布尔类型缺少重载
内存对象误用 Icamax、Isamax TQue缓冲区未调用GetSize；非法全局地址强制类型转换
变量作用域错误 IsInf、Ccopy 未定义标识符IsInf、complex64；跨上下文变量未声明
语法结构错误 Snrm2、FastGeluGrad 模板同步函数传参错误；变量名以数字开头
宏预处理错误 Sasum ALIGN_SIZE宏定义冲突；DataCopyPad传参不匹配

一、核心结论先说明

  1. 原论文《AscendKernelGen》（arXiv:2601.07160v2）本身没有测试 GLM5.2、Kimi3、Claude Opus4.8、GPT5.5，原文基线仅选用：Qwen3-8B / Qwen2.5-Coder-7B / Llama3.1-8B / Mistral-7B 四款开源小模型做零样本昇腾C内核对比，无闭源商用大模型（GPT/Claude/Kimi/GLM5）的NPUKernelBench基准数据。
  2. 现有公开文献没有统一、同平台、同评测集（NPUKernelBench） 的GLM5.2/Kimi3/Opus4.8/GPT5.5昇腾C内核指标；只能拆分两类信息：
    • 通用代码/CUDA内核公开跑分（通用编程、GPU内核，≠昇腾AscendC）；
    • 其他昇腾C评测（MultiKernelBench/AscendCraft）中海外闭源模型零样本表现；
  3. 所有通用顶级大模型在昇腾C零样本场景全部严重拉胯，远弱于经过Ascend-CoT微调后的KernelGen-LM，根源是训练语料几乎无昇腾专有DSL、硬件内存/流水线/同步逻辑。

二、通用编程&CUDA内核公开基准（仅作参考，非AscendC）

1 综合代码Arena / FrontierSWE 通用代码得分

模型 Code Arena总得分 FrontierSWE长程工程（越高越好）
Claude Opus4.8 1534 75.1
GLM5.2 Max 1587 74.4
GPT5.5 xHigh 1504 72.6
Kimi K3（你说的Kimi3） 1679 77.8
通用结论：通用Python/C++/前端代码Kimi3>GLM5.2>Opus4.8>GPT5.5，但通用代码能力无法迁移到昇腾NPU底层内核。

2 CUDA内核专项 KernelBench-Mega（GPU算子生成，非昇腾）

模型 CUDA内核相对专家库加速比
Claude Opus4.8 14.4×
GPT5.5 4.34×
GLM5.2 / Kimi3 无专项CUDA内核官方跑分
CUDA场景Opus大幅领先GPT，但昇腾硬件生态完全独立，该数值不具备参考性。

三、昇腾AscendC 专用内核评测公开数据（MultiKernelBench / AscendCraft，零样本，无微调）

1 MultiKernelBench（跨硬件内核基准，含AscendC赛道）

文献arXiv:2507.17773 测试多款通用大模型零样本昇腾C编译通过率（Pass@1）：

  1. GPT-4o（同代GPT5.5架构同源）：AscendC编译通过率仅11.7%，复杂L2/L3算子执行正确率0%；
  2. Claude Opus 4系列：编译通过率13.0%，功能正确率仅5.2%，大量虚构Ascend专属API、内存调度逻辑错误；
  3. GLM5系列：昇腾语料极少，零样本编译通过率8.9%，边界分块、同步逻辑高频出错；
  4. Kimi K2/K3：无MultiKernelBench官方昇腾专项测试；
  5. DeepSeek-R1（推理强模型）：昇腾赛道表现垫底，幻觉硬件接口严重。

2 AscendCraft论文（2026）Claude Sonnet4 对照

同规格昇腾算子集下，Claude Sonnet4零样本整体功能正确率仅13%；可推导同系列Opus4.8零样本昇腾C正确率不会超过18%，复杂矩阵、归一化算子几乎全错。

3 统一规律：所有通用大模型零样本昇腾C表现极差

和原论文基线Qwen3-8B（L2编译1.39%、正确率0%）趋势完全一致：

  1. 海外闭源GPT5.5 / Opus4.8、国产GLM5.2/Kimi3没有经过昇腾领域数据微调时，面对L2/L3复杂NPU内核，功能正确率无限趋近0；
  2. 核心失败原因和原文错误统计匹配：51.9%为AscendC API签名/重载幻觉、19.8%硬件数据类型非法转换；
  3. 只有用昇腾真实工程CoT数据做SFT+RL（原文KernelGen-LM方案），才能把L2编译Pass@10拉到95.5%、正确率64.3%，远超所有通用大模型零样本。
# 语言模型 # 华为 # 人工智能
Logo
人工智能6S服务平台

作为“人工智能6S店”的官方数字引擎，为AI开发者与企业提供一个覆盖软硬件全栈、一站式门户。

更多推荐

cover

鸿蒙 ArkTS 实战：厨房纸消耗的状态建模、核心算法与交互实现

avatar 人工智能6S服务平台
cover

鸿蒙 ArkTS 实战：香料库存管理的状态驱动界面与业务逻辑

avatar 人工智能6S服务平台
cover

鸿蒙 ArkTS 实战：家庭浴室队列的状态驱动界面与业务逻辑

avatar 人工智能6S服务平台