Ascend C资料贡献指南

概述

Ascend C资料体系包含五份核心文档，开发者可以通过提交PR对资料进行改进和贡献：

文档	内容	文档目录
入门教程	Ascend C概述、环境准备、快速上手（HelloWorld、首个算子）	`docs/zh/guide/getting_started`
编程指南	编程模型、编程范式、编译运行、硬件架构、高级编程	`docs/zh/guide/programming_guide`
API参考手册	接口参数、约束、示例、API关联	`docs/zh/api/`
算子实践参考	算子实现、性能优化、调优案例	`docs/zh/guide/operator_practice/`
跨代迁移兼容性指南	API兼容策略、架构变更映射、迁移步骤	`docs/zh/guide/cross_gen_migration_guide/`

在阅读本文档前，请确保您已了解昇腾AI处理器硬件架构。零基础开发者推荐先阅读入门教程；已有基础的开发者推荐先阅读Ascend C编程指南。

贡献场景

文档纠错

如果您在资料中发现描述错误、参数值不准确、约束遗漏等问题：

按照社区指引新建 Documentation | 文档反馈 类Issue，指出对应文档的问题
在评论框中输入 /assign 或 /assign @yourself 将该Issue分配给您
修复后提交PR

文档补充

如果您发现某些内容在资料中缺失或不够完整（如缺少示例、缺少约束说明、缺少关联API说明）：

新建 Requirement | 需求建议 类Issue，描述需要补充的内容
按照本指南的"编写规范"完成补充
提交PR

性能优化案例贡献

如果您有Ascend C算子性能优化的实战经验愿意分享：

参考 operator_practice/best_practices/ 目录下已有的调优案例结构
按照本指南的"性能优化案例编写规范"编写文档
提交PR

编写规范

Ascend C资料遵循《Ascend C资料设计规范》的三个维度要求：可获取性、可读性、完备性。以下是针对常见编写场景的具体规范。

通用规范

规则	要求	错误示例	正确示例
术语分层	高层章节用抽象名，低层章节首次出现硬件名时关联高层名	概述文档直接写"数据通过UB搬入VEC"	概述文档写"数据从外部存储搬入片上存储"，低层文档写"数据从Global Memory搬入UB（片上存储）"
易混概念区分	名称相近的概念必须用对比表区分，不能靠读者自行推断	只说"四步法"不区分 Tiling流程和TPipe流水	用表格区分：四步法=编程流程，TPipe四步=流水管理范式
约束前置醒目	硬件约束在概念首次出现时标注，不放在后续章节或仅靠编译报错暴露	32B对齐约束只在assert中出现	在参数说明前用独立段落标注"约束与限制"
首次引入加链接	提到其他文档负责的内容时，首次出现加链接，后续不重复	每次提到DataCopy都加链接	仅首次提到DataCopy时加链接到API参考手册
不逆序阅读	概念B依赖概念A时，A先出现或B处有明确引用链接	API页面直接用TPipe概念但未铺垫	API页面开头放"前置知识"段，链接到编程指南对应章节

入门教程编写规范

入门教程是零基础入口，面向初次接触Ascend C的开发者，帮助其快速建立全景认知并完成首个算子。

文件职责：

文件类型	职责	禁止
概述与学习路径	全景介绍：Ascend C是什么、学习路径推荐	深入技术细节（链接到编程指南）
环境准备	实操步骤：安装、配置、验证环境	重复编程指南中的编译运行详解
快速入门/异构系统与编程模型	入门概念：Host/Device、AI Core、SIMD/SIMT选型	深入编程模型原理（链接到编程指南）
快速入门/基于SIMD编程	实践：HelloWorld + 首个算子（Add）	重复编程指南的完整编程范式
快速入门/基于SIMT编程	实践：HelloWorld + 首个算子（Gather）	重复编程指南的完整编程范式

目录结构映射（与编程指南章节对应，不含技术附录）：

入门教程目录	对应编程指南章节	入门教程定位
getting_started/	编程模型/编程模型概述	概述与学习路径
environment_setup.md	编译与运行	快速环境搭建，详细编译说明链接到编程指南
quick_start/heterogeneous_system_and_programming_model.md	编程模型/异构系统 + 编程模型/编程模型概述	入门级概念+SIMD/SIMT选型，深入链接到编程指南
quick_start/simd_programming/	编程模型/AI-Core-SIMD编程	HelloWorld+Add算子快速上手
quick_start/simt_programming/	编程模型/AI-Core-SIMT编程	HelloWorld+Gather算子快速上手

链接方向：

首次提到编程概念深入 → 链接到编程指南对应章节
首次提到某个API名称 → 链接到API参考手册
不需要链接到算子实践参考或跨代迁移指南（入门阶段不涉及）

编程指南编写规范

编程指南是概念权威源，其他文档遇到编程概念问题都链接回编程指南。

文件职责：

文件类型	职责	禁止
概述/总览文件	导航页：列出子主题+一句话概述+链接	展开技术细节
概念介绍文件	定义概念、解释原理、给出约束	重复其他文件内容（改为引用）
操作指南文件	代码片段+操作步骤+注意事项	重复概念定义（链接到概念文件）

编程指南到其他文档的链接方向：

首次提到某个API名称 → 链接到API参考手册
首次提到优化/实践话题 → 链接到算子实践参考
提到架构版本差异 → 链接到跨代迁移指南

API参考页面编写规范

API参考手册是接口详情权威源，每个API页面必须包含以下要素（按顺序）：

页面结构总览（按出现顺序）

序号	节标题	是否必须	说明
1	产品支持情况	✅ 必须	页面最顶部，用列表形式列出每个产品系列的支持状态
2	功能说明	✅ 必须	头文件路径 + 总述段落 + 数学公式/示意图（3要素必须齐全）
3	矩阵/张量计算说明表	条件必须	矩阵计算类API必须有，按产品系列分表
4	函数原型	✅ 必须	所有重载原型，每个原型一个代码块
5	参数说明	✅ 必须	主参数表 + 嵌套结构体参数表
6	数据类型	条件必须	矩阵/向量计算类API必须；工具类/配置类API不需要
7	返回值说明	条件必须	有返回值的API必须有
8	约束说明	✅ 必须	单参数约束就近说明；多参数/组合约束独立段落描述
9	关键特性说明	条件必须	矩阵/向量计算类API如涉及关键硬件特性（HF32、GEMV、UnitFlag等），须有对应特性说明节；简单API可省略
10	理论性能（附录引用）	条件推荐	不再在API页面内展开cycle计算公式，改为引用附录中的理论性能汇总表
11	调用示例	✅ 必须	代码片段示例 + 链接到样例库（如有对应样例，如果接口没有对应样例可省略链接）
12	关联API	条件必须	同一版本内的互斥API + 功能相近API

1. 产品支持情况

格式要求：用列表形式列出每个产品系列的支持状态。

样板写法：

## 产品支持情况

- Ascend 950PR/Ascend 950DT：支持
- Atlas A3 训练系列产品/Atlas A3 推理系列产品：支持
- Atlas A2 训练系列产品/Atlas A2 推理系列产品：支持
- Atlas 200I/500 A2 推理产品：不支持

编写要求：

每个产品系列一行，格式为 - 产品名：支持/不支持
部分支持标注为 支持* 并在后续说明限制条件

2. 功能说明

格式要求：3要素必须齐全——头文件路径 + 总述段落 + 数学公式/示意图。头文件路径放在功能说明最开头。

样板写法（以Mmad为例）：

## 功能说明

头文件为: `#include "basic_api/kernel_operator_mm_intf.h"`

Mmad接口是AscendC面向昇腾AI芯片的矩阵乘加核心计算接口，专为高性能算子开发设计，
封装了昇腾NPU硬件的矩阵乘加计算能力...

如下图所示，Mmad接口实现昇腾NPU矩阵乘计算能力，其数学表达式为：
C = A × B + Bias
[示意图/公式图片]

编写要求：

头文件必须写明具体include路径，放在功能说明最开头（不独立成节，作为功能说明的一部分）
总述段说明API的定位（面向什么场景、封装了什么硬件能力）
数学公式必须给出（不能只用文字描述）
如有示意图/数据流图，嵌入图片并标注图序号

3. 矩阵/张量计算说明表

适用范围：矩阵计算类API（Mmad、LoadData、Fixpipe等）必须有此表。非矩阵计算API跳过。

格式要求：按产品系列分表，列出参与计算的张量的物理位置、维度、格式。

样板写法（以Mmad为例）：

针对Atlas A2/A3系列产品：
| 矩阵计算 | 物理位置 | 维度 | 输入/输出数据格式 | 数据类型 |
|----------|----------|------|------------------|----------|
| A | L0A Buffer | M x K | Zz | 数据类型 |
| B | L0B Buffer | K x N | Zn | 数据类型 |
| C | L0C Buffer | M x N | Nz | 数据类型 |

针对Ascend 950PR/Ascend 950DT产品：
| 矩阵计算 | 物理位置 | 维度 | 输入/输出数据格式 | 数据类型 |
|----------|----------|------|------------------|----------|
| A | L0A Buffer | M x K | Nz | 数据类型 |
| B | L0B Buffer | K x N | Zn | 数据类型 |
| C | L0C Buffer | M x N | Nz | 数据类型 |

编写要求：

当不同产品的排布格式不同（如Zz vs Nz）时，必须分表
物理位置写Buffer名称（L0A/L0B/L0C/Unified Buffer（UB）等）
格式写分形排布格式（Zz/Zn/Nz/ND等）
如有偏置矩阵，在C矩阵行的说明中标注"可支持使用偏置矩阵Bias进行初始化，维度为1 x N"

4. 函数原型

格式要求：按功能分支分组列出所有重载。每个原型单独一个代码块，配1-2行功能说明。

样板写法（以Mmad为例）：

// 不支持bias：基础矩阵乘加原型
template <typename T, typename U, typename S>
__aicore__ inline void Mmad(const LocalTensor<T>& c, const LocalTensor<U>& a,
    const LocalTensor<S>& b, const MmadParams& mmadParams)

// 支持bias：带偏置矩阵乘加原型
template <typename T, typename U, typename S, typename V>
__aicore__ inline void Mmad(const LocalTensor<T>& c, const LocalTensor<U>& a,
    const LocalTensor<S>& b, const LocalTensor<V>& bias, const MmadParams& mmadParams)

// 不传入bias（BitMode版本）：位域联合体参数原型
template <typename T, typename U, typename S>
__aicore__ inline void Mmad(const LocalTensor<T>& c, const LocalTensor<U>& a,
    const LocalTensor<S>& b, const MmadBitModeParams& mmadParams)

// 传入bias（BitMode版本）：位域联合体+偏置原型
template <typename T, typename U, typename S, typename V>
__aicore__ inline void Mmad(const LocalTensor<T>& c, const LocalTensor<U>& a,
    const LocalTensor<S>& b, const LocalTensor<V>& bias, const MmadBitModeParams& mmadParams)

编写要求：

每个模板参数（T, U, S, V...）必须在参数说明中标注其对应关系
__aicore__ 修饰符必须保留
参数按输入顺序列出：先输出（c），再输入（a, b, bias），最后参数结构体
如有BitMode版本（位域联合体参数），与常规版本并列列出

5. 参数说明

格式要求：分两层——主参数表 + 嵌套结构体参数表。

主参数表

| 参数名称 | 输入/输出 | 含义 |
|----------|----------|------|
| c | 输出 | 目的操作数，结果矩阵c，类型为LocalTensor，物理存储位置为L0C Buffer(TPosition:CO1)。起始地址需按1024字节对齐。 |
| a | 输入 | 源操作数，左矩阵a，类型为LocalTensor，物理存储位置为L0A Buffer(TPosition:A2)。起始地址需按512字节对齐。 |
| b | 输入 | 源操作数，右矩阵b，类型为LocalTensor，物理存储位置为L0B Buffer(TPosition:B2)。起始地址需按512字节对齐。 |
| bias | 输入 | 源操作数，bias矩阵，物理存储位置为BT Buffer(TPosition:C2)。起始地址需按64字节对齐。 |
| mmadParams | 输入 | 矩阵乘相关参数。MmadParams参数说明请参考下表。 |

编写要求：

含义列必须包含：功能描述 + LocalTensor的物理位置(TPosition枚举值) + 对齐字节数
对齐要求写入含义列（不单独列行），格式："起始地址需要按照N字节对齐"
复杂参数（结构体/联合体）含义列末尾加"详细说明请参考下表"

嵌套结构体参数表

### MmadParams结构体内参数说明

| 参数名称 | 含义 |
|----------|------|
| m | 左矩阵Height，取值范围：m∈[0, 4095]。默认值为0。 |
| n | 右矩阵Width，取值范围：n∈[0, 4095]。默认值为0。 |
| k | 左矩阵Width/右矩阵Height，取值范围：k∈[0, 4095]。默认值为0。 |
| cmatrixInitVal | 是否使能C矩阵默认初始化清零。true=初始化为0，false=不初始化（由cmatrixSource控制）。默认true。 |
| cmatrixSource | C矩阵初始值是否来源于BT Buffer。true=来源BT，false=不初始化。默认false。**注意**：带bias输入时此参数无效。 |
| unitFlag | 控制Mmad和Fixpipe细粒度并行。0=不使能，2=使能不复位，3=使能并复位。详见UnitFlag特性说明。 |
| disableGemv | M=1时是否关闭GEMV模式。**仅Ascend 950PR/Ascend 950DT支持**。 |

编写要求：

每个参数写明取值范围（数学区间）和默认值
参数如有产品差异，在该参数行内直接标注"仅xxx产品支持"或分列写不同产品取值
参数间的互斥/依赖关系用注意块标注（如"带bias时此参数无效"）
废弃参数标注 > **已废弃** 并给出替代方案
BitMode参数类（union+bit-field设计），需额外说明设计思想，列出每个bit位的含义和Get/Set函数

6. 数据类型

格式要求：分产品系列列出类型组合表。以Mmad API为例，分"无bias"和"有bias"两张表。

样板写法（以Mmad为例）：

针对Atlas A2/A3系列产品：

| 左矩阵A | 右矩阵B | 结果矩阵C |
|---------|---------|-----------|
| int8_t | int8_t | int32_t |
| half | half | float |
| float | float | float |
| bfloat16_t | bfloat16_t | float |
| int4b_t | int4b_t | int32_t |

针对Ascend 950PR/Ascend 950DT产品：
（类似分表，含fp8/hifloat8等特有类型）

| 左矩阵A | 右矩阵B | 偏置Bias | 结果矩阵C |
|---------|---------|---------|-----------|
| int8_t | int8_t | int32_t | int32_t |
| half | half | float | float |
| ... | ... | ... | ... |

编写要求：

必须分产品系列——不同芯片支持的类型组合不同（如950PR新增fp8/hifloat8）
无bias和有bias是不同的表（bias列有额外的类型约束）
表头用类型名（不用缩写），按精度从低到高排列
每种组合是一行，不能合并写"int8_t/int4b_t"

7. 返回值说明（条件必须）

适用范围：有返回值的API（如Cast返回LocalTensor引用、GetBlockIdx返回uint32_t等）。

格式要求：说明返回值类型和含义。

样板写法：

## 返回值说明

| 返回值类型 | 含义 |
|----------|------|
| LocalTensor<T>& | 返回目的操作数的引用，类型与模板参数T一致。可用于链式调用。 |

编写要求：

表格列出返回值类型和含义
如果返回值可链式调用，标注"可用于链式调用"

8. 约束说明

格式要求：分两层描述。单个参数的约束在参数说明表格内就近说明（如对齐要求、取值范围）；多参数间的约束关系、多接口组合使用的约束，在独立的"约束说明"段落中集中描述。

样板写法：

## 约束说明

- 结果矩阵C只支持位于L0C Buffer（CO1），左矩阵A只支持位于L0A Buffer（A2），右矩阵B只支持位于L0B Buffer（B2）。
- 当M、K、N中的任意一个值为0时，指令不执行，视为NOP（空操作）。**应尽量避免NOP.M操作**，NOP.M实测远大于1 cycle。
- 当M=1时，默认开启GEMV模式，A矩阵按ND格式读取（不视为ZZ/NZ格式），起始地址仍要求512字节对齐。
- 一次Mmad至少完成A(16×16×half)×B(16×16×half)的数据块计算。有效值非16倍数时无效数据排布如下：
  [示意图]

### 同步优化说明
当矩阵沿K轴累加时，连续两次Mmad间是否需要PipeBarrier(PIPE_M)取决于计算量：

if ((m / 16) * (n / 16) < 10) {
    AscendC::PipeBarrier<PIPE_M>();  // 计算量小于阈值，需同步
}
// 计算量大于阈值时，硬件自动处理依赖，无需同步

### UnitFlag特性约束
Mmad和Fixpipe的unitFlag需同步开启。前n-1条设为2（维持占用），最后一条设为3（解除占用）。
[详细说明见UnitFlag特性说明](链接)

### 特殊值/边界值约束
- 浮点INF/NAN：通过CTRL[48]设置饱和/非饱和模式
- **Mmad应避免NAN输入**，否则可能产生执行报错
- 整数类型只有饱和模式

编写要求：

位置约束必须优先描述物理存储位置，后面加括号标注TPosition枚举值
空操作/NOP约束必须写明避免场景和性能影响
分形粒度必须写明最小计算块大小和无效数据排布规则
同步约束给出明确的阈值公式和代码片段
特殊值处理给出CTRL寄存器设置方法
每个复杂特性（UnitFlag/kDirectionAlign/GEMV等）给出独立小节 + 链接到详细特性说明

9. 关键特性说明（条件必须）

适用范围：矩阵/向量计算类API如涉及关键硬件特性（如HF32模式、GEMV加速、UnitFlag并行控制等），须有此节。简单API（如配置类、工具类）可省略。

格式要求：按特性分小节，每个特性用标题+表格/段落说明。 样板写法（以Mmad为例）：

## 关键特性说明

### HF32模式
通过SetHF32Mode/SetHF32TransMode配置HF32（Hybrid Float32）模式，实现float类型矩阵乘以9bit指数共享的方式减少计算量。
| 模式 | 说明 |
|------|------|
| HF32_OP0 | 指数共享左矩阵A的列方向 |
| HF32_OP1 | 指数共享右矩阵B的行方向 |

### GEMV加速
当M=1时，Mmad自动触发GEMV（矩阵向量乘）加速通路，跳过不必要的矩阵加载。

### UnitFlag并行控制
控制Mmad和Fixpipe的细粒度流水并行，详见[UnitFlag特性说明](api/SIMD-API/basic_api/cube_compute_ISASI/mmad_compute_key_features/UnitFlag.md)。

### K方向对齐约束
k方向需满足特定的对齐要求，不同数据类型约束不同，详见[K方向对齐约束](api/SIMD-API/basic_api/cube_compute_ISASI/mmad_compute_key_features/k_direction_alignment_constraint.md)。

编写要求：

每个特性独立小节，标题使用###级标题
特性说明包含：是什么、什么时候用、怎么配置、有什么约束
复杂特性（如HF32涉及多个配置接口联动）用表格对比不同模式
每个特性节末尾链接到详细特性说明文档（如有独立文档）

10. 理论性能（附录引用）

格式要求：不在API页面内展开cycle计算公式和并行度参数表，改为引用附录中的理论性能汇总。

样板写法：

## 理论性能

本接口的理论性能数据请参考附录：
- Cube指令理论性能汇总：详见[Cube指令理论性能汇总](api/appendix/cube_instruction_theoretical_perf_summary.md)
- Vector指令理论性能汇总：详见[Vector指令理论性能汇总](api/appendix/vector_instruction_theoretical_perf_summary.md)

编写要求：

理论性能数据统一维护在附录中，API页面不再内联cycle计算公式
引用方式为链接到对应附录页（Cube类API引用Cube附录，Vector类API引用Vector附录）
如API有特殊的性能注意事项（如首指令开销、特定参数下的性能差异），可在引用附录后补充说明

11. 调用示例

格式要求：链接到样例库 + 场景分类表。

样板写法：

## 调用示例

### 示例代码
```cpp
#include "kernel_operator.h"

extern "C" __global__ __aicore__ void mmad_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    // 算子内部实现逻辑...
}

矩阵乘的完整样例请参考[Mmad样例](链接)。

编写要求：

链接到样例仓库的具体目录（如有对应的样例）。如果接口当前没有对应样例，可省略样例链接。

12. 关联API（条件必须）

格式要求：分三个层次——互斥API（警告框）→ 相似API选择指引（表格）→ 底层映射（折叠）。

互斥API

用 > **警告** 框醒目标注不能同时使用或功能冲突的API。

> **警告**：以下接口与当前接口互斥，不能在同一Tiling策略中同时使用：
> - **MmadWithSparse**：当启用结构化稀疏（4选2）时使用，与Mmad互斥。
> - **MmadMx**：MX矩阵计算场景使用，与Mmad互斥。

> **注意**：Mmad的带bias原型和不带bias原型**不互斥**，可根据场景自由选择。

相似API选择指引

用表格列出功能相近的API，给出选择建议，帮助开发者决策。

| 相似API | 与当前API区别 | 选择建议 |
|---------|-------------|---------|
| [MmadWithSparse](api/SIMD-API/basic_api/cube_compute_ISASI/mmad_compute/MmadWithSparse.md) | 支持4选2结构化稀疏矩阵乘 | 权重矩阵已剪枝为稀疏格式时选用 |
| [MmadMx](api/SIMD-API/basic_api/cube_compute_ISASI/mmad_compute/MmadMx.md) | 支持MX格式（微缩放）矩阵乘 | 使用MX混合精度训练时选用 |
| [Iterate](api/SIMD-API/adv_api/cube_compute/Matmul_Kernel/Iterate.md) + [GetTensorC](api/SIMD-API/adv_api/cube_compute/Matmul_Kernel/GetTensorC.md) | Matmul高阶API的循环迭代模式 | 不需要精细控制L0A/L0B排布时，优先用高阶API |
| Mmad带bias原型 | 支持偏置矩阵初始化C | 需要C+=A×B+Bias场景选用，比cmatrixSource配置更简单 |
| Mmad不带bias + cmatrixSource | 通过参数配置C初始值来源 | 需要C初始值来自BT Buffer但不需要Bias时选用 |
| Mmad + MmadBitModeParams | 位域联合体参数，单入参传递 | 追求极致性能、需要细粒度bit位操作时选用 |

底层built-in映射（可选，折叠展示）

列出高级API到硬件built-in指令的封装关系，帮助深度调优开发者理解底层参数映射。

<details>
<summary>底层built-in接口映射（点击展开）</summary>

Mmad接口在built-in接口（mad）基础上抽象封装：

针对Atlas A2/A3：
void mad(__cc__ float *c, __ca__ half *a, __cb__ half *b,
    uint16_t m, uint16_t k, uint16_t n, uint8_t unitFlag,
    bool kDirectionAlign, bool cmatrixSource, bool cmatrixInitVal);

针对Ascend 950PR/Ascend 950DT：
void mad(__cc__ float *c, __ca__ half *a, __cb__ half *b,
    uint16_t m, uint16_t k, uint16_t n, uint8_t unit_Flag_ctrl,
    bool gemv_ctrl, bool BTbuf_ctrl, bool zero_Cmatrix_ctrl);
</details>

算子实践参考编写规范

算子实践参考是实践案例源，负责展开编程指南点到为止的实践和优化内容。

文档结构模板：

# <算子名> 算子实践参考

## 前置说明
阅读本文档前，您需要了解：xxx概念（链接到编程指南）、xxx API（链接到API参考）。

## 算子实现
### 基础版本
（代码片段 + 说明 + 链接到样例库完整示例）

### 进阶版本
（如适用：性能优化版本、多数据类型版本等）

## 性能优化
（如适用：列出针对此算子的优化手段，链接到对应优化专题）

## 常见问题
（如适用：开发过程中容易踩的坑）

关键要求：

有多种实现方案（如MemBase vs RegBase）的算子，开头放方案差异表+选择建议
代码中首次出现的API加链接到API参考手册，后续不重复
优化方案如果因架构版本不同，标注 [适用版本：仅xxx] 并链接到跨代迁移指南

跨代迁移页面编写规范

跨代迁移指南是兼容性权威源，负责说明架构版本间的差异和迁移路径。

迁移映射条目格式：

## <遗留API名> → <新API名>

### 变更说明
（一句话说明为什么变更、什么时候变更的）

### 参数差异对照

| 参数 | 旧接口 | 新接口 | 差异说明 |
|------|--------|--------|---------|
| ... | ... | ... | ... |

### 约束差异
（列出对齐、数据类型、元素数量等方面的差异）

### 迁移代码示例
（展示旧写法和新写法的对比）

### 迁移验证
（迁移后如何验证功能和性能）

关键要求：

每个映射条目必须链接到新旧API的API参考手册详情页
概念变更必须链接回编程指南对应章节
使用统一的 [适用版本：xxx] 标注格式

代码示例编写规范

代码片段要求

要求	说明
精简聚焦	只展示当前讲解相关的代码，用 `// ... 其他初始化` 省略无关部分
可理解	每个片段配2-3行文字说明这段代码做了什么
可追踪	片段末尾链接到样例库完整示例
关键字注释	Ascend C特有关键字（`__aicore__`、`__ubuf__`、`__simd_vf__`等）必须行内注释
参数覆盖	覆盖常用参数组合，不仅展示单一用法

示例中的关键字注释

__aicore__ inline void ExampleKernel(__gm__ uint8_t* x) {  // __aicore__=核函数（Kernel）修饰符, __gm__=Global Memory地址空间
    // ...
}

性能优化案例编写规范

参考 operator_practice/best_practices/ 目录下的已有案例（如FlashAttention、Matmul系列）。推荐结构：

# <算子名> 性能调优案例

## 问题背景
（此算子在实际场景中的性能瓶颈是什么）

## 优化思路
（从哪些维度优化：Tiling、内存访问、流水编排、计算指令选择等）

## 优化实现
### 基线版本
（未优化的代码片段）

### 优化版本
（优化后的代码片段，逐项说明每处改动的原因）

## 性能对比

| 指标 | 基线 | 优化后 | 提升幅度 |
|------|------|--------|---------|
| 带宽利用率 | ... | ... | ... |
| 计算吞吐 | ... | ... | ... |

## 适用范围
[适用版本：xxx]（如果优化方案因架构版本不同）

提交PR前的自检清单

提交资料PR前，请逐项检查：

内容正确性：

所有硬件参数值（容量、粒度、范围）与硬件规格保持一致
"固定值"等绝对性描述已经核实（很多"固定值"实际由配置参数决定）
数据通路描述与实际路径一致（不与其他文档矛盾）

约束完备性：

每个API的支持数据类型列表完整
地址对齐要求已明确（字节数和方向）
元素数量范围（最小值、最大值、对齐粒度）已说明
API组合限制（互斥、必需搭配）已标注
跨版本差异已用 [适用版本：xxx] 标注

链接规范：

首次提到其他文档负责的API/概念时已加链接
同一概念/API后续重复出现不重复加链接
API页面的"前置知识"段已链接到编程指南
涉及版本差异时已链接到跨代迁移指南

可读性：

术语使用与章节抽象层级匹配（高层用抽象名，低层首次出现时关联）
易混概念用对比表区分（不用纯文字暗示）
代码片段精简聚焦，有文字说明，有样例库链接
Ascend C特有关键字有行内注释

示例质量：

每个API至少有1个最小可运行示例
示例可复制编译运行
覆盖多数据类型/多参数组合（不只有单一用法）

资料体系与链接关系

五份文档之间通过交叉链接形成导航网络，遵循"谁提到别人负责的内容，谁就加链接"原则：

入门教程 ──编程概念深入──→ 编程指南
         ──首次引入API──→ API参考手册

编程指南 ──首次引入API──→ API参考手册
        ──首次引入优化──→ 算子实践参考
        ──版本差异────→ 跨代迁移指南

API参考手册 ──前置概念──→ 编程指南
            ──版本差异──→ 跨代迁移指南

算子实践参考 ──首次使用API──→ API参考手册
            ──编程概念────→ 编程指南
            ──跨架构优化──→ 跨代迁移指南

跨代迁移指南 ──概念定义──→ 编程指南
            ──新API详情──→ API参考手册

样例库（asc-devkit/examples/）不属于资料体系，但五份文档中的代码示例均可链接到样例库。

Ascend C资料贡献指南

概述

贡献场景

文档纠错

文档补充

性能优化案例贡献

编写规范

通用规范

入门教程编写规范

编程指南编写规范

API参考页面编写规范

页面结构总览（按出现顺序）

1. 产品支持情况

2. 功能说明

3. 矩阵/张量计算说明表

4. 函数原型

5. 参数说明

主参数表

嵌套结构体参数表

6. 数据类型

7. 返回值说明（条件必须）

8. 约束说明

9. 关键特性说明（条件必须）

10. 理论性能（附录引用）

11. 调用示例

12. 关联API（条件必须）

互斥API

相似API选择指引

底层built-in映射（可选，折叠展示）

算子实践参考编写规范

跨代迁移页面编写规范

代码示例编写规范

代码片段要求

示例中的关键字注释

性能优化案例编写规范

提交PR前的自检清单

资料体系与链接关系

更多信息