• 人工智能
• 算子库
• 深度学习
• CANN
• Ascend

AddRmsNormQuantV2 是 CANN ops-nn 算子库中面向大模型部署的融合算子，它将"Add 残差相加 → RmsNorm 归一化 → 双路静态量化"三段计算收敛为单算子执行，相比 V1 版本新增了偏置项与归一化中间结果输出能力。本文将以其官方文档为骨架，结合仓库中的算子定义、ACLNN 两段式接口实现、kernel 源码与单测用例，系统讲解该算子的功能语义、参数约束、产品差异、调用流程与底层实现路径，帮助开发者在 norm/add_rms_norm_quant_v2 模块内快速上手并正确使用。

本文以 norm/add_rms_norm_quant_v2/README.md 与 norm/add_rms_norm_quant_v2/docs/aclnnAddRmsNormQuantV2.md 为权威依据，源码佐证统一指向当前仓库路径。

一、为什么需要这个算子：从三段式到单算子融合

在大模型推理与训练中，残差连接（x = x1 + x2）之后紧跟 RmsNorm 归一化是 Transformer 类网络的标准结构。RmsNorm 相比 LayerNorm 去掉了"减去均值"的步骤，只对输入做均方根归一化，因此更适合长序列场景下的大规模并行计算。

如果使用原始算子组合，数据需要经历"Add 写回 → RmsNorm 读入/写出 → Quantize 读入量化"多次完整的搬入搬出（GM↔L1/L0 往返）。AddRmsNormQuantV2 将三段计算融合为一个 kernel，从算子层面消除了中间结果的反复搬运；在此基础上，V2 版本相较于 AddRmsNormQuant 增加了两个关键能力：

1. 偏置项bias（betaOptional）：RmsNorm 归一化后叠加可选的逐元素偏置；
2. 可选输出resOut（rmsNormOut）：直接输出不含 bias 的归一化结果，供上层网络复用，避免重复计算。

其计算语义在 op_graph/add_rms_norm_quant_v2_proto.h 中有完整的伪代码注释：

x = x1 + x2 rstd = np.rsqrt(np.mean(np.power(x, 2), reduce_axis, keepdims=True) + epsilon) resOut = x * rstd * gamma rmsnorm_out = resOut + bias // div_mode 为 True 时 y1 = round(rmsnorm_out / scales1 + zero_points1) y2 = round(rmsnorm_out / scales2 + zero_points2) // div_mode 为 False 时 y1 = round(rmsnorm_out * scales1 + zero_points1) y2 = round(rmsnorm_out * scales2 + zero_points2)

二、算子功能与计算公式

2.1 三段计算链路

AddRmsNormQuantV2 一次完成如下三个阶段（对应公式符号与文档一致）：

阶段一：残差相加

$$x_i = {x1}_i + {x2}_i$$

阶段二：RmsNorm 归一化并叠加偏置

$$y_i=\frac{1}{\operatorname{Rms}(\mathbf{x})} \cdot x_i \cdot \gamma_i + \text{bias}i,\quad \operatorname{Rms}(\mathbf{x})=\sqrt{\frac{1}{n}\sum{i=1}^{n} x_i^2 + \epsilon}$$

阶段三：双路静态量化（round 到目标量化类型）

• div_mode = True（除法模式，默认）：

$$y1_i = round(y_i / scales1_i + zero_points1_i),\quad y2_i = round(y_i / scales2_i + zero_points2_i)$$

• div_mode = False（乘法模式）：

$$y1_i = round(y_i \cdot scales1_i + zero_points1_i),\quad y2_i = round(y_i \cdot scales2_i + zero_points2_i)$$

可选归一化中间结果（不含 bias）：

$$resOut_i = \frac{1}{\operatorname{Rms}(\mathbf{x})} \cdot x_i \cdot \gamma_i$$

值得注意：bias与resOut是相互独立的，resOut的计算不包含 bias 项，二者可以任意组合使用（有或无 bias、输出或不输出 resOut 均可），这一设计在 docs/aclnnAddRmsNormQuantV2.md 的"约束说明"中有明确说明。

2.2 产品支持情况

从 op_host/add_rms_norm_quant_v2_def.cpp 的算子注册可以看到对应内核配置：ascend910b、ascend910_93、ascend310p以及ascend950（其中ascend910_93对应 Atlas A3 系列）。ascend950配置声明了DynamicCompileStaticFlag(true).DynamicRankSupportFlag(true).DynamicShapeSupportFlag(true)，即支持动态 shape 编译。

三、参数说明与产品差异

3.1 输入 / 输出 / 属性总表

以下为参数表中的数据类型与数据格式并集（"√"表示必选或支持，详见各产品差异说明）：

属性默认值在算子定义文件 op_host/add_rms_norm_quant_v2_def.cpp 与 IR 定义 op_graph/add_rms_norm_quant_v2_proto.h 中均有完全一致的注册：axis=-1、epsilon=1e-6f、div_mode=true、dst_type=DT_INT8、output_res=false。

3.2 产品差异说明（数据类型与属性）

参数表是全部支持产品能力的并集，实际使用必须按产品对齐。约定："x类参数"包括x1、x2、gamma、bias、x、resOut；"scale参数"包括scales1和非空scales2；"zero point参数"包括非空zero_points1、zero_points2。同一合法组合内同类参数数据类型保持一致。

关键回退机制：对 Atlas A3、A2 与推理系列，通过 aclnn API 调用且div_mode=False时，接口不进入原生 V2 路径；当输出x、不输出resOut且其他参数满足 V1 回退约束时，接口回退到 V1 路径（即 AddRmsNormQuant）并按乘法模式执行。GE 图模式不适用该回退机制，div_mode仅支持 True。

这一机制在 op_host/op_api/aclnn_add_rms_norm_quant_v2.cpp 的CheckSupportV2中落地：当传入scales2Optional、zeroPoints2Optional或divMode=false时返回 false（不支持原生 V2），随后在aclnnAddRmsNormQuantV2GetWorkspaceSize主流程（同文件 L876-L885）中判断：CheckSupportV2为 true 则走ComputeAddRmsNormQuantV2，否则若xOut非空则回退到ComputeAddRmsNormQuantV1（内部调用 l0op 的AddRmsNormQuant并做 ViewCopy 拼接）。

3.3 产品差异说明（shape 与参数组合）

其中"归一化元素个数不小于 32"在 API 层由 aclnn_add_rms_norm_quant_v2.cpp 的CheckGammaShapeX落地：310P 平台上当gamma各维乘积（即归一化元素总数）小于SIZE_MIN_LIMIT_310P = 32时直接返回ACLNN_ERR_PARAM_INVALID。

3.4 约束说明汇总

• 可选输出x与resOut：必须且只能选择其一进行输出。
• 需要输出y2时：gamma与scales1的 shape 保持一致，且与x1需要 Norm 的维度保持一致，此时可选输出只能输出x。
• 输出x时：若scales1、zero_points1的 shape 为[1]，且gamma为 1 维（大小等于x1最后一维）或 2 维（第一维为 1、第二维等于x1最后一维），则scales2、zero_points2不生效。
• 输出resOut时：scales1、zero_points1的 shape 为[1]，gamma满足上述 1 维/2 维约束，且bias与x必须传空指针，此时scales2、zero_points2不生效。
• 维度边界：x1、x2、gamma、bias、scales1、scales2、zero_points1、zero_points2、y1、y2、x、resOut的 shape 每一维大小均不大于 INT32 最大值 2147483647。
• 数据格式：所有输入输出 Tensor 推荐使用 ND 格式，其他格式会由框架默认转换成 ND 处理。

四、aclnn 两段式接口与调用流程

4.1 函数原型

每个算子都遵循 CANN 单算子调用的两段式接口规范：先调用GetWorkspaceSize接口获取 workspace 大小与执行器，再调用执行接口完成计算。

aclnnStatus aclnnAddRmsNormQuantV2GetWorkspaceSize( const aclTensor *x1, const aclTensor *x2, const aclTensor *gamma, const aclTensor *scales1, const aclTensor *scales2Optional, const aclTensor *zeroPoints1Optional, const aclTensor *zeroPoints2Optional, const aclTensor *betaOptional, int64_t axis, double epsilon, bool divMode, aclTensor *y1Out, aclTensor *y2Out, aclTensor *xOut, aclTensor *rmsNormOut, uint64_t *workspaceSize, aclOpExecutor **executor); aclnnStatus aclnnAddRmsNormQuantV2( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream);

对应头文件为 op_host/op_api/aclnn_add_rms_norm_quant_v2.h，调用时包含#include "aclnnop/aclnn_add_rms_norm_quant_v2.h"。

4.2 第一段接口参数要点

• x1、x2、gamma、scales1、y1Out、y2Out为必选；scales2Optional、zeroPoints1Optional、zeroPoints2Optional、betaOptional为可选，均支持传入空指针（表示"不使用"）。
• x2：shape 与数据类型需与x1保持一致。
• gamma：shape 与x1需要归一化的维度保持一致，数据类型与x1一致。
• scales2Optional：shape、数据类型与scales1一致；divMode=True时值不能为 0。
• zeroPoints1Optional：shape 与scales1一致。
• zeroPoints2Optional：shape 与scales1一致，数据类型与zeroPoints1Optional一致。
• betaOptional：shape 与数据类型需与gamma一致。
• axis：int64，仅支持 -1（需小于等于x1维度数），传其他值不生效。
• epsilon：double，建议传较小的正数，防止除 0。
• divMode：bool，True 用除法、False 用乘法。
• 所有 tensor 支持 1-8 维、支持非连续 Tensor（框架内部会做 contiguous 转换），数据格式 ND。

4.3 返回值与错误码

aclnnStatus返回状态码约定见 docs/zh/context/aclnn_return_code.md。第一段接口完成入参校验，常见报错：

五、完整调用示例与关键步骤拆解

仓库提供了可直接编译运行的样例 examples/test_aclnn_add_rms_norm_quant_v2.cpp（编译与运行方式参考 docs/zh/context/compile_and_run_sample.md）。下面按步骤拆解其核心逻辑。

步骤 1：初始化 device 与 stream

auto ret = aclInit(nullptr); ret = aclrtSetDevice(deviceId); ret = aclrtCreateStream(&stream);

步骤 2：构造输入输出 aclTensor

样例以xShape = {64, 32}、gammaShape = {32}为例，即x1为 64×32 的 FLOAT16 张量，gamma为长度 32 的 1 维权重，与x1的最后一维对应。通过aclrtMalloc申请 device 内存、aclrtMemcpy拷入数据、aclCreateTensor创建 ND 格式连续 tensor：

std::vector<int64_t> xShape = {64, 32}; std::vector<int64_t> gammaShape = {32}; std::vector<int16_t> x1HostData(xShapeSize, 0); // FLOAT16 std::vector<float> scales1HostData(gammaShapeSize, 1); std::vector<int32_t> zeroPoints1HostData(gammaShapeSize, 100); float epsilon = 1e-6; int64_t axis = -1; bool divMode = true;

步骤 3：第一段接口（计算 workspace 并建执行器）

注意可选参数位置：样例传入了beta，而scales2、zeroPoints2、rmsNormOut传空指针：

uint64_t workspaceSize = 0; aclOpExecutor* executor; ret = aclnnAddRmsNormQuantV2GetWorkspaceSize( x1, x2, gamma, scales1, nullptr /*scales2*/, zeroPoints1, nullptr /*zeroPoints2*/, beta, axis, epsilon, divMode, y1, y2, x, rmsNormOut, &workspaceSize, &executor);

步骤 4：申请 workspace 并执行第二段接口

if (workspaceSize > 0) { aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); } ret = aclnnAddRmsNormQuantV2(workspaceAddr, workspaceSize, executor, stream);

步骤 5：同步并取回结果

ret = aclrtSynchronizeStream(stream); ret = aclrtMemcpy(resultData.data(), size * sizeof(int8_t), y1DeviceAddr, size * sizeof(int8_t), ACL_MEMCPY_DEVICE_TO_HOST);

步骤 6-7：资源释放

依次aclDestroyTensor释放各 aclTensor，aclrtFree释放 device 内存（含 workspace），最后aclrtDestroyStream、aclrtResetDevice、aclFinalize。

样例使用的数据类型组合（x 类 = FLOAT16，scales = FLOAT32，zero points = INT32，输出 = INT8）对应 README 中"FLOAT16/FLOAT32/INT32"这一合法组合，也是 op_host/config/ascend950/add_rms_norm_quant_v2_binary.json 中注册的 binary 组合之一。

六、GE 图模式调用方式

除 aclnn API 外，AddRmsNormQuantV2 还支持通过 GE 图模式构图调用，IR 定义位于 op_graph/add_rms_norm_quant_v2_proto.h，使用REG_OP(AddRmsNormQuantV2)注册算子：

• 输入：x1、x2、gamma、scales1（必选），scales2、zero_points1、zero_points2、bias（可选）；
• 输出：y1、y2（必选），x、resOut（可选）；
• 属性：axis（默认 -1）、epsilon（默认 1e-6）、div_mode（默认 true）、dst_type（默认 DT_INT8）、output_res（默认 false）。

注意 GE 图模式没有 aclnn 的 V1 回退机制，div_mode仅支持 True。

七、底层实现：kernel 分派与动态编译

kernel 入口位于 op_kernel/add_rms_norm_quant_v2.cpp，其实现复用了 V1 算子的三类 kernel 模板，通过 tiling key 区分不同计算模式：

extern "C" __global__ __aicore__ void add_rms_norm_quant_v2(...) { if (TILING_KEY_IS(10)) KernelAddRmsNormQuant<...> op(&pipe); else if (TILING_KEY_IS(10010)) KernelAddRmsNormQuantSplitD<...> op(&pipe); else if (TILING_KEY_IS(30010)) KernelAddRmsNormQuantSingleN<...> op(&pipe); // 101 / 10101 / 30101 / 1011 / 11011 / 31011 等分支…… }

从源码结构看，该算子针对不同 shape 切分策略（普通 / SplitD / SingleN）与不同模式组合分别编译了多种 kernel 变体。模板参数false, true, false、true, false, true等差异对应不同的功能开关（如是否有 bias、是否输出 resOut 等计算分支），实际由 host 侧 tiling 数据在编译时决定。kernel 内核对x1/x2/gamma/scales1/scales2/zero_points1/zero_points2/bias做op.Init(...)初始化后调用op.Process()执行完整计算流水。

此外，Ascend 950 路径使用独立实现 op_kernel/add_rms_norm_quant_v2_apt.cpp，并在 op_host/add_rms_norm_quant_v2_def.cpp 中通过ExtendCfgInfo("opFile.value", "add_rms_norm_quant_v2_apt")指定，配合 op_host/config/ascend950/add_rms_norm_quant_v2_binary.json 中的 binary 清单按 dtype 组合（fp16/fp32/int32、bf16/fp32/int32、hifloat8、float8_e5m2、float8_e4m3fn 等）做静态编译分发。

八、单元测试与确定性

仓库在 tests/ut/op_host/test_aclnn_add_rms_norm_quant_v2.cpp 中提供了基于 GTest 的 aclnn 接口测试，覆盖 Ascend 950PR 的静态（_stc_）与动态（_dyn_）shape 用例，例如：

TEST_F(l2_add_rms_norm_quant_v2_test, ascend950PR_9589_case_dyn_001) { auto tensor_desc_x1 = TensorDesc({8, 64}, ACL_FLOAT16, ACL_FORMAT_ND); auto tensor_desc_x2 = TensorDesc({8, 64}, ACL_FLOAT16, ACL_FORMAT_ND); auto tensor_desc_gamma = TensorDesc({64}, ACL_FLOAT16, ACL_FORMAT_ND); auto tensor_desc_y1 = TensorDesc({8, 64}, ACL_INT8, ACL_FORMAT_ND); // ... auto ut = OP_API_UT(aclnnAddRmsNormQuantV2, INPUT(tensor_desc_x1, tensor_desc_x2, tensor_desc_gamma, tensor_desc_s1, tensor_desc_s2, (aclTensor*)nullptr, (aclTensor*)nullptr, tensor_desc_beta, axis, eps, divMode), OUTPUT(tensor_desc_y1, tensor_desc_y2, tensor_desc_x, (aclTensor*)nullptr)); uint64_t workspace_size = 0; aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size); }

同一目录下还有 kernel 层测试 tests/ut/op_kernel/test_add_rms_norm_quant_v2.cpp 与对应的 def 头文件。此外文档明确 aclnnAddRmsNormQuantV2 为默认确定性实现，适合在需要稳定复现结果的训练/推理链路中使用。

九、实践要点速查

1. 先确认产品型号：Ascend 950PR&950DT 能力最全（支持 FP8 输出与div_mode=False）；Atlas A3/A2 与推理系列输出仅 INT8，原生 V2 仅支持div_mode=True。
2. 数据类型保持同类一致：同一合法组合内，x 类参数之间、scale 参数之间、zero point 参数之间、y1/y2之间类型必须一致；950 上scales1=FLOAT32时zero_points只能是 FLOAT32 或 INT32。
3. 正确理解可选输出：x与resOut二选一；输出resOut时bias与x必须传空指针；输出y2时只能带x。
4. div_mode=True时 scales 不能为 0；Atlas A3/A2 与推理系列上zero_points2非空时scales2必须同时非空。
5. 310P 平台：归一化元素个数（gamma各维乘积）不得小于 32，输入不支持 inf/NaN。
6. 回退机制：Atlas 系列通过 aclnn 调用且div_mode=False时，满足条件会自动回退到 V1 乘法路径；GE 图模式无此回退。
7. 推荐 ND 格式，其余格式由框架转换；每维大小不超过 INT32 最大值。

十、相关文档索引

• 算子 README：norm/add_rms_norm_quant_v2/README.md
• aclnn 接口文档：norm/add_rms_norm_quant_v2/docs/aclnnAddRmsNormQuantV2.md
• 两段式接口说明：docs/zh/context/two_phase_api.md
• aclnn 返回码说明：docs/zh/context/aclnn_return_code.md
• 编译与运行样例：docs/zh/context/compile_and_run_sample.md
• 算子 IR 定义：op_graph/add_rms_norm_quant_v2_proto.h
• 算子定义注册：op_host/add_rms_norm_quant_v2_def.cpp
• aclnn 实现：op_host/op_api/aclnn_add_rms_norm_quant_v2.cpp
• 调用样例：examples/test_aclnn_add_rms_norm_quant_v2.cpp
• kernel 实现：op_kernel/add_rms_norm_quant_v2.cpp
• 950 binary 配置：op_host/config/ascend950/add_rms_norm_quant_v2_binary.json
• 单测用例：tests/ut/op_host/test_aclnn_add_rms_norm_quant_v2.cpp

• 人工智能
• 算子库
• 深度学习
• CANN
• Ascend