人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载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 增加了两个关键能力偏置项biasbetaOptionalRmsNorm 归一化后叠加可选的逐元素偏置可选输出resOutrmsNormOut直接输出不含 bias 的归一化结果供上层网络复用避免重复计算。其计算语义在 op_graph/add_rms_norm_quant_v2_proto.h 中有完整的伪代码注释x x1 x2 rstd np.rsqrt(np.mean(np.power(x, 2), reduce_axis, keepdimsTrue) 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{i1}^{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 产品支持情况产品是否支持Ascend 950PR 950DT 系列产品√Atlas A3 系列产品√Atlas A2 系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品√Atlas 训练系列产品×从 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 输入 / 输出 / 属性总表以下为参数表中的数据类型与数据格式并集√表示必选或支持详见各产品差异说明参数名输入/输出/属性描述数据类型数据格式x1输入源数据张量对应公式x1FLOAT16、BFLOAT16、FLOAT32NDx2输入源数据张量对应公式x2shape 与x1一致FLOAT16、BFLOAT16、FLOAT32NDgamma输入归一化权重张量对应公式gammashape 与x1需要 Norm 的维度一致FLOAT16、BFLOAT16、FLOAT32NDscales1输入第一路量化 scales对应公式scales1div_modeTrue时不能为 0FLOAT32、FLOAT16、BFLOAT16NDscales2可选输入第二路量化 scales对应公式scales2可不传div_modeTrue时不能为 0shape 与scales1一致FLOAT32、FLOAT16、BFLOAT16NDzero_points1可选输入第一路量化 offset对应公式zero_points1可不传不叠加 offsetshape 与gamma一致INT32、FLOAT32、FLOAT16、BFLOAT16NDzero_points2可选输入第二路量化 offset对应公式zero_points2可不传shape 与gamma一致与scales2的合法组合见产品差异INT32、FLOAT32、FLOAT16、BFLOAT16NDbias可选输入归一化偏置项对应公式bias可不传不叠加偏置shape 与gamma一致FLOAT16、BFLOAT16、FLOAT32NDy1输出第一路量化输出对应公式y1shape 与x1/x2一致INT8、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FNNDy2输出第二路量化输出对应公式y2shape 与x1/x2一致有效条件见产品差异INT8、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FNNDx可选输出x1与x2之和对应公式xshape 与输入一致FLOAT16、BFLOAT16、FLOAT32NDresOut可选输出不含bias的 RmsNorm 结果对应公式resOutoutput_resTrue时输出FLOAT16、BFLOAT16、FLOAT32NDaxis可选属性量化 element-wise 轴其他轴广播当前仅支持 -1传入其他值不生效默认 -1INT64-epsilon可选属性防止除 0 错误对应公式epsilon默认 1e-6FLOAT32-div_mode可选属性量化使用除法还是乘法默认 TrueBOOL-dst_type可选属性y1/y2的数据类型默认 DT_INT8INT64-output_res可选属性是否输出不含bias的resOut默认 FalseBOOL-属性默认值在算子定义文件 op_host/add_rms_norm_quant_v2_def.cpp 与 IR 定义 op_graph/add_rms_norm_quant_v2_proto.h 中均有完全一致的注册axis-1、epsilon1e-6f、div_modetrue、dst_typeDT_INT8、output_resfalse。3.2 产品差异说明数据类型与属性参数表是全部支持产品能力的并集实际使用必须按产品对齐。约定x类参数包括x1、x2、gamma、bias、x、resOutscale参数包括scales1和非空scales2zero point参数包括非空zero_points1、zero_points2。同一合法组合内同类参数数据类型保持一致。产品x类/scale/zero point 合法组合y1、y2数据类型dst_typediv_modeAscend 950PR950DTFLOAT16/FLOAT32/INT32、BFLOAT16/FLOAT32/INT32、FLOAT32/FLOAT32/FLOAT32、FLOAT16/FLOAT16/FLOAT16、BFLOAT16/BFLOAT16/BFLOAT16、FLOAT16/FLOAT32/FLOAT32、BFLOAT16/FLOAT32/FLOAT32INT8、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FNy1与y2类型一致True、FalseAtlas A3 / Atlas A2FLOAT16/FLOAT32/INT32、BFLOAT16/BFLOAT16/BFLOAT16、BFLOAT16/FLOAT32/INT32INT8原生 V2 路径仅 Trueaclnn API 满足 V1 回退条件时支持 FalseAtlas 推理系列FLOAT16/FLOAT32/INT32INT8原生 V2 路径仅 Trueaclnn API 满足 V1 回退条件时支持 False关键回退机制对 Atlas A3、A2 与推理系列通过 aclnn API 调用且div_modeFalse时接口不进入原生 V2 路径当输出x、不输出resOut且其他参数满足 V1 回退约束时接口回退到 V1 路径即 AddRmsNormQuant并按乘法模式执行。GE 图模式不适用该回退机制div_mode仅支持 True。这一机制在 op_host/op_api/aclnn_add_rms_norm_quant_v2.cpp 的CheckSupportV2中落地当传入scales2Optional、zeroPoints2Optional或divModefalse时返回 false不支持原生 V2随后在aclnnAddRmsNormQuantV2GetWorkspaceSize主流程同文件 L876-L885中判断CheckSupportV2为 true 则走ComputeAddRmsNormQuantV2否则若xOut非空则回退到ComputeAddRmsNormQuantV1内部调用 l0op 的AddRmsNormQuant并做 ViewCopy 拼接。3.3 产品差异说明shape 与参数组合产品静态/动态 shape第二路量化其他用户可观察限制Ascend 950PR950DT输入、输出均支持 NDscales2或zero_points2任一非空时y2有效仅zero_points2非空时scales2按 1 处理$y2round(yzero_points2)$-Atlas A3 / Atlas A2输入、输出均支持 NDy2仅在scales2非空时有效zero_points2非空时scales2必须同时非空输入为 inf 时输出 inf输入为 NaN 时输出 NaNAtlas 推理系列输入、输出均支持 NDy2仅在scales2非空时有效zero_points2非空时scales2必须同时非空输入不支持 inf 和 NaNgamma包含的归一化元素个数不能小于 32其中归一化元素个数不小于 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为可选均支持传入空指针表示不使用。x2shape 与数据类型需与x1保持一致。gammashape 与x1需要归一化的维度保持一致数据类型与x1一致。scales2Optionalshape、数据类型与scales1一致divModeTrue时值不能为 0。zeroPoints1Optionalshape 与scales1一致。zeroPoints2Optionalshape 与scales1一致数据类型与zeroPoints1Optional一致。betaOptionalshape 与数据类型需与gamma一致。axisint64仅支持 -1需小于等于x1维度数传其他值不生效。epsilondouble建议传较小的正数防止除 0。divModeboolTrue 用除法、False 用乘法。所有 tensor 支持 1-8 维、支持非连续 Tensor框架内部会做 contiguous 转换数据格式 ND。4.3 返回值与错误码aclnnStatus返回状态码约定见 docs/zh/context/aclnn_return_code.md。第一段接口完成入参校验常见报错返回码错误码描述ACLNN_ERR_PARAM_NULLPTR161001必选输入、输出或必选属性为空指针ACLNN_ERR_PARAM_INVALID161002输入或输出的数据类型不在支持范围内五、完整调用示例与关键步骤拆解仓库提供了可直接编译运行的样例 examples/test_aclnn_add_rms_norm_quant_v2.cpp编译与运行方式参考 docs/zh/context/compile_and_run_sample.md。下面按步骤拆解其核心逻辑。步骤 1初始化 device 与 streamauto 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 格式连续 tensorstd::vectorint64_t xShape {64, 32}; std::vectorint64_t gammaShape {32}; std::vectorint16_t x1HostData(xShapeSize, 0); // FLOAT16 std::vectorfloat scales1HostData(gammaShapeSize, 1); std::vectorint32_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释放各 aclTensoraclrtFree释放 device 内存含 workspace最后aclrtDestroyStream、aclrtResetDevice、aclFinalize。样例使用的数据类型组合x 类 FLOAT16scales FLOAT32zero 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 为默认确定性实现适合在需要稳定复现结果的训练/推理链路中使用。九、实践要点速查先确认产品型号Ascend 950PR950DT 能力最全支持 FP8 输出与div_modeFalseAtlas A3/A2 与推理系列输出仅 INT8原生 V2 仅支持div_modeTrue。数据类型保持同类一致同一合法组合内x 类参数之间、scale 参数之间、zero point 参数之间、y1/y2之间类型必须一致950 上scales1FLOAT32时zero_points只能是 FLOAT32 或 INT32。正确理解可选输出x与resOut二选一输出resOut时bias与x必须传空指针输出y2时只能带x。div_modeTrue时 scales 不能为 0Atlas A3/A2 与推理系列上zero_points2非空时scales2必须同时非空。310P 平台归一化元素个数gamma各维乘积不得小于 32输入不支持 inf/NaN。回退机制Atlas 系列通过 aclnn 调用且div_modeFalse时满足条件会自动回退到 V1 乘法路径GE 图模式无此回退。推荐 ND 格式其余格式由框架转换每维大小不超过 INT32 最大值。十、相关文档索引算子 READMEnorm/add_rms_norm_quant_v2/README.mdaclnn 接口文档norm/add_rms_norm_quant_v2/docs/aclnnAddRmsNormQuantV2.md两段式接口说明docs/zh/context/two_phase_api.mdaclnn 返回码说明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.cppaclnn 实现op_host/op_api/aclnn_add_rms_norm_quant_v2.cpp调用样例examples/test_aclnn_add_rms_norm_quant_v2.cppkernel 实现op_kernel/add_rms_norm_quant_v2.cpp950 binary 配置op_host/config/ascend950/add_rms_norm_quant_v2_binary.json单测用例tests/ut/op_host/test_aclnn_add_rms_norm_quant_v2.cpp赞分享人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载相关推荐CANN ops-nn 的 AddRmsNormQuant 算子Add RmsNorm 双路 Quantize 融合的标准化量化实现解析CANN ops nn 的 AddRmsNormQuant 算子Add RmsNorm 双路 Quantize 融合的标准化量化实现解析 导读 Add人工智能算子库深度学习CANNAscendDeepSeek 提示词输入框接入 Vim 模式编辑Voyager vimInput 声明式插件原理与验证DeepSeek 提示词输入框接入 Vim 模式编辑Voyager vimInput 声明式插件原理与验证 Voyager 是一个面向 Gemini、AI S人工智能算子库深度学习CANNAscendTiptap 2.x 到 3.x 升级5 步完成富文本编辑器迁移Tiptap 2.x 到 3.x 升级5 步完成富文本编辑器迁移 本文面向仍停留在 2.x 的 Tiptap 使用者讲清 Tiptap 3.x 迁移要动哪些人工智能算子库深度学习CANNAscend上一篇train-llm-from-scratch 优化与训练系统实战从 AdamW、余弦学习率到梯度累积与 DDP 多卡训练下一篇Haystack 中的 LibreOfficeFileConverter用 soffice 打通遗留 Office 文档的格式转换管线创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考