算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载本文以 CANN 开源数学算子库 ops-math 中 math/bias_add_grad/README.md 为骨架结合该算子目录下的算子定义、形状推导、Tiling、核函数与调用示例源码系统讲解 BiasAddGrad 算子的功能语义、参数约束、产品支持情况以及如何在 NPU 上通过算子 IR 构图方式完成图模式调用。读完本文你将能看懂 BiasAddGrad 算子的完整工程实现链路并可直接参考仓库内的调用样例在自己的推理/训练图接入该算子。一、算子功能与数学原理BiasAddGrad 是BiasAdd偏置加法的反向算子用于计算偏置bias的梯度。在前向网络中BiasAdd将一维偏置向量加到输入张量上反向传播时则需要将输出梯度沿除通道channel/feature维度以外的所有维度进行累加还原出偏置向量的梯度供后续优化器更新偏置参数使用。其核心计算公式如下来自 README$$ \text{y}[c] \sum_{n0}^{N-1} \sum_{h0}^{H-1} \sum_{w0}^{W-1} \text{x}[n, c, h, w] $$即输入张量 x 的通道维下标 c 固定时其余所有维度batch 维 n、空间维 h/w 等的元素全部累加得到输出 y 中下标为 c 的一个标量。因此输出 y 是一维张量其长度等于输入张量的通道数feature 维大小。从算子原型注释可以进一步确认其语义见 op_graph/bias_add_grad_proto.hPerforms the backward operation for BiasAdd on the bias tensor. It accumulates all the values from out_backprop into the feature dimension. For NHWC data format, the feature dimension is the last. For NCHW data format, the feature dimension is the third-to-last.同时该文件还明确说明其兼容性与 TensorFlow 的 BiasAddGrad 算子兼容Compatible with the TensorFlow operator BiasAddGrad并通过 framework/bias_add_grad_tf_plugin.cpp 注册为FrameworkType(TENSORFLOW)、OriginOpType(BiasAddGrad)的自定义算子映射从而支持 TensorFlow 图在昇腾 NPU 上的迁移。二、产品支持情况根据 math/bias_add_grad/README.md 的产品支持情况章节该算子支持的产品如下产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品√Atlas 训练系列产品√从源码层面看该支持范围与算子编译配置一致CMakeLists.txt 中set(SUPPORT_COMPUTE_UNIT ascend950 ascend350)并在 op_host/bias_add_grad_def.cpp 中通过this-AICore().AddConfig(ascend950, aicoreConfig)和AddConfig(ascend350, aicoreConfig)为两类计算单元注册 AICore 配置同时 op_host/config 目录下分别维护了ascend950与ascend350两套算子二进制描述bias_add_grad_binary.json和 tiling key 配置bias_add_grad_simplified_key.ini。三、参数说明算子包含两个张量参数一个输入、一个输出与一个可选属性参数明细如下表来自 README 参数说明章节参数名输入/输出/属性描述数据类型数据格式x输入待进行 BiasAdd 计算的入参公式中的 xBFLOAT16、FLOAT16、FLOAT、INT8、INT32、INT64、UINT8NDy输出累加偏差公式中的 y_cBFLOAT16、FLOAT16、FLOAT、INT8、INT32、INT64、UINT8NDdata_format属性表示输入张量的格式可选STRING-3.1 关于 data_format 属性data_format是一个字符串属性用于指定输入张量的排布方式从而决定哪个维度是特征维通道维、即累加目标维。在算子原型中该属性有明确默认值op_graph/bias_add_grad_proto.hREG_OP(BiasAddGrad) .INPUT(x, TensorType::NumberType()) .OUTPUT(y, TensorType::NumberType()) .ATTR(data_format, String, NHWC) .OP_END_FACTORY_REG(BiasAddGrad)默认值为NHWC此时特征维是张量的最后一维下标dim_num - 1取值为NCHW时特征维是第 1 维下标 1即从 0 数起的第二个维度。该取值约束在形状推导实现中被显式校验仅接受NCHW与NHWC两种取值见 op_host/bias_add_grad_infershape.cppstatic const std::vectorstd::pairstd::string, Format kFormatMap { {NCHW, FORMAT_NCHW}, {NHWC, FORMAT_NHWC}, }; // data_format 不在 (NCHW, NHWC) 中时直接返回 GRAPH_FAILED OP_CHECK_IF(it kFormatMap.end(), OP_LOGE(context-GetNodeName(), data_format %s must be in (NCHW, NHWC)., data_format), return GRAPH_FAILED);3.2 数据类型实现的进一步说明README 参数表列出了较宽的数据类型范围BFLOAT16、FLOAT16、FLOAT、INT8、INT32、INT64、UINT8。从仓库源码看算子定义阶段实际注册的 AICore 数据类型为FLOAT16、FLOAT、BFLOAT16、INT32、INT64五种op_host/bias_add_grad_def.cppstatic const std::vectorge::DataType dataType {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT64};与之一致的是 op_host/config/ascend950/bias_add_grad_binary.json 中按 dtype 分文件编译的算子二进制清单BiasAddGrad_float16、BiasAddGrad_float32、BiasAddGrad_bfloat16、BiasAddGrad_int32、BiasAddGrad_int64输入输出均采用ND格式format_match_mode为FormatAgnosticshape 为[-2]表示支持动态 shape。在具体接入前建议以目标环境的算子支持矩阵为准进行 dtype 确认。四、约束说明README 中约束说明章节明确标注为无约束。结合源码可以进一步理解为该算子在工程层面做了充分的动态化支持从而降低了使用约束具体体现在 op_host/bias_add_grad_def.cpp 的 AICore 配置中aicoreConfig.DynamicCompileStaticFlag(true) .DynamicRankSupportFlag(true) .DynamicShapeSupportFlag(true) .ExtendCfgInfo(opFile.value, bias_add_grad_apt);DynamicCompileStaticFlag(true)支持动态编译与静态编译两种模式DynamicRankSupportFlag(true)支持动态 rank输入维度数量可变DynamicShapeSupportFlag(true)支持动态 shapeExtendCfgInfo(opFile.value, bias_add_grad_apt)指明核函数源文件为bias_add_grad_apt对应 op_kernel/bias_add_grad_apt.cpp。唯一需要注意的约束是NCHW 格式下输入张量的维度数必须大于等于 2否则形状推导会报错见 op_host/bias_add_grad_infershape.cppif (dim_num DIM_SIZE2) { OP_LOGE(context-GetNodeName(), The bias add grad op dimension(%lu) must be greater than or equal to 2 when format is NCHW!, dim_num); return GRAPH_FAILED; }五、输出形状与数据类型推导5.1 形状推导输出 y 的形状是一维的其唯一维度取自输入 x 的特征维大小推导逻辑位于 op_host/bias_add_grad_infershape.cppout_shape-SetDimNum(0); out_shape-AppendDim(in_shape-GetDim((it-second FORMAT_NHWC) ? (dim_num - 1) : 1));即data_format NHWC时输出维度 in_shape[dim_num - 1]最后一维data_format NCHW时输出维度 in_shape[1]第二维。例如输入 shape 为[3, 4]、格式为 NHWC见调用样例输出 y 的 shape 即为[4]。5.2 数据类型推导输出 y 与输入 x 保持相同的数据类型该规则通过图推导接口声明op_graph/bias_add_grad_graph_infer.cppIMPL_OP(BiasAddGrad).InferDataType(InferDataTypeOutputSameAsInput);六、图模式调用实践README 的调用说明章节给出了算子的标准调用方式图模式调用。仓库提供了完整的可运行示例 examples/test_geir_bias_add_grad.cpp通过算子 IRop_graph/bias_add_grad_proto.h 构图的方式调用 BiasAddGrad 算子。整体流程可拆解为以下五个阶段6.1 初始化 GE 环境示例首先调用ge::GEInitialize完成 GEGraph Engine全局初始化并通过全局选项指定执行设备与运行模式test_geir_bias_add_grad.cppstd::mapAscendString, AscendString global_options {{ge.exec.deviceId, 0}, {ge.graphRunMode, 1}}; Status ret ge::GEInitialize(global_options);其中ge.exec.deviceId指定使用的 NPU 设备号ge.graphRunMode指定图运行模式。6.2 构图创建 BiasAddGrad 算子节点构图核心在CreateOppInGraph函数中test_geir_bias_add_grad.cpp。示例先通过op::BiasAddGrad(bias_add_grad)创建算子节点再借助文件顶部的ADD_INPUT/ADD_OUTPUT/ADD_INPUT_ATTR宏完成输入、输出与属性设置auto add1 op::BiasAddGrad(bias_add_grad); vectorvectorint64_t shapes {{3, 4}, {3, 4}}; ADD_INPUT(1, x, inDtype, shapes[0]); // 输入 xshape [3, 4] ADD_OUTPUT(2, y, inDtype, shapes[1]); // 输出 yshape [3, 4] ADD_INPUT_ATTR(data_format, NHWC); // 属性 data_format NHWC需要注意示例中输出 y 的 shape 按[3, 4]构造只是占位描述实际运行时的输出形状由前一节的形状推导规则确定NHWC 下取最后一维即[4]。宏内部还会调用GenOnesData生成全 2 的输入数据并依次完成TensorDesc构造FORMAT_ND 输入 dtype、graph.AddOp添加算子节点、set_input_x建立数据依赖等动作test_geir_bias_add_grad.cpp。6.3 创建会话并添加图构图完成后示例创建ge::Session会话并通过AddGraph将计算图注册到会话中test_geir_bias_add_grad.cppge::Session* session new Session(build_options); uint32_t graph_id 0; Status ret session-AddGraph(graph_id, graph, graph_options);6.4 执行图并导出输入输出数据会话就绪后调用aclgrphDumpGraph将图 dump 为 txt 文件用于调试再通过session-RunGraph(graph_id, input, output)执行算子test_geir_bias_add_grad.cppstd::string file_path ./dump; aclgrphDumpGraph(graph, file_path.c_str(), file_path.length()); Status ret session-RunGraph(graph_id, input, output);执行结果会被写入本地二进制文件tc_ge_irrun_test_0008_npu_input_i.bin与tc_ge_irrun_test_0008_npu_output_i.bin并逐元素打印输出结果test_geir_bias_add_grad.cpp便于与预期结果比对。示例中还有GetDataTypeSize工具函数test_geir_bias_add_grad.cpp按 DataType 换算单元素字节数供数据落盘时计算缓冲区大小。6.5 资源清理最后调用ge::GEFinalize完成 GE 反初始化并释放会话资源同时通过GEGetErrorMsgV2/GEGetWarningMsgV2输出运行期的错误与告警信息test_geir_bias_add_grad.cpp。七、核函数与调度实现BiasAddGrad 的 NPU 侧计算由基于模板化算子开发框架atvoss实现的核函数完成体现出累加归约这一核心运算本质。7.1 核函数入口核函数定义在 op_kernel/bias_add_grad_apt.cpptemplate REDUCE_TPL_PARAM __global__ __aicore__ void bias_add_grad(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { if (g_coreType AIC) { return; } ... REGISTER_TILING_DEFAULT(ReduceOpTilingData); GET_TILING_DATA_WITH_STRUCT(ReduceOpTilingData, tilingData, tiling); TPipe pipe; using PromoteType __reduceType::GetPromoteTypeDTYPE_X::T; using Op ReduceSchREDUCE_TPL_VALUE, BiasAddGrad::BiasAddGradDagDTYPE_X, PromoteType::OpDag; Op op(tilingData); op.Init(pipe, x, y, userWS); op.Process(); }可以看到该算子复用了 atvoss 的归约调度器ReduceSch并通过GetPromoteType为不同类型选取累加时的提升类型PromoteType保证低精度累加过程的精度。7.2 计算 DAG算子内部的计算流水线以模板 DAG 形式描述见 op_kernel/bias_add_grad_dag.husing OpCopyIn0 BindVec::CopyInT, Placeholder::In0T; // 输入拷入 using Cast0 BindVec::CastPromteT, T, 0, OpCopyIn0; // 提升精度 using ReduceOp0 BindVec::ReduceSumOpPromteT, Cast0; // 归约求和 using Cast1 BindVec::CastT, PromteT, 1, ReduceOp0; // 恢复原精度 using OpCopyOut BindVec::CopyOutT, Placeholder::Out0T, Cast1; // 结果写出 ... using MemCfg MemOptCfgMemLevel::LEVEL_2; // 使用 L2 做归约内存优化该 DAG 清晰对应了本章开头给出的数学公式CopyIn → 提升类型 Cast → ReduceSum → 恢复类型 Cast → CopyOut即按通道维做全局求和同时在MemLevel::LEVEL_2上做内存优化配置。7.3 Tiling 与模板实例化模板实例化策略op_kernel/bias_add_grad_tiling_key.h 通过ASCENDC_TPL_SEL按归约模式空模式、A、AR、ARA、ARAR 及各自的 Group 变体声明算子模板参数组合用于在编译期展开不同数据流形态的实例Host 侧 Tilingop_host/arch35 目录下提供bias_add_grad_tiling.cpp/.h与bias_add_grad_tiling_arch35.cpp/.h其中 bias_add_grad_tiling.h 定义BiasAddGradCompileInfo结构体内部封装 atvoss 的ReduceOpCompileInfo即复用统一的归约 Tiling 框架计算分块参数芯片适配CMakeLists.txt 中set(SUPPORT_TILING_DIR arch35 arch35)表示 ascend950 与 ascend350 均复用 arch35 目录下的 Tiling 实现。八、测试与验证仓库为 BiasAddGrad 提供了 Host 侧单测与 ST 用例便于开发者验证实现正确性形状推导单测tests/ut/op_host/test_bias_add_grad_infershape.cpp 验证 NCHW/NHWC 两种格式下的输出形状推导是否符合输出长度 通道维大小的规则Tiling 单测tests/ut/op_host/arch35/test_bias_add_grad_tiling.cpp 验证 Host 侧 Tiling 参数计算ST 用例tests/st/arch35/ttk_kernel_bias_add_grad_st.csv 描述内核级 ST 的输入输出规格结果比对脚本tests/assets/golden.py 提供期望值golden生成逻辑供自动化测试比对 NPU 实际输出。九、小结BiasAddGrad 是 BiasAdd 的反向算子核心语义为沿通道维归约求和输出一维偏置梯度。在 CANN ops-math 仓库中它具备完整的工程链路通过 算子原型 声明接口与默认属性通过 算子定义 注册 AICore 配置支持 ascend950/ascend350、动态 rank/shape通过 形状推导 决定输出维度并通过基于 atvoss 归约框架的 核函数 与 计算 DAG 完成 NPU 侧累加。若要快速上手可直接参考 图模式调用示例按照初始化 GE → 构图 → 建会话 → 执行 → 清理的标准流程接入该算子。赞分享算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载相关推荐BatchToSpaceND 算子深度解析原理、约束与 NPU 图模式调用实践CANN ops-mathBatchToSpaceND 算子深度解析原理、约束与 NPU 图模式调用实践CANN ops math BatchToSpaceND 是 CANN op算子库人工智能CANNJarkViewer看图软件从零开始完整上手3分钟装好AVIF、HEIC、实况照片全能看JarkViewer看图软件从零开始完整上手3分钟装好AVIF、HEIC、实况照片全能看 你是不是也遇到过这样的场景手机里刚拍的实况照片、相机导出的RAW算子库人工智能CANNCANN ops-math 算子 MatrixDiagPartV3 深度解析对角线提取算法、参数配置与图模式调用实践CANN ops math 算子 MatrixDiagPartV3 深度解析对角线提取算法、参数配置与图模式调用实践 MatrixDiagPartV3 是 C算子库人工智能CANN上一篇uBlock Origin 使用指南5 分钟装好广告拦截器并验证生效下一篇如何在Java应用中完美渲染LaTeX数学公式JLaTeXMath终极解决方案指南创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
