算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载RealDiv 是 CANN ops-math 数学算子库中实现逐元素实数除法y x1 / x2的基础算子覆盖从 Ascend 950 到 Atlas A2/A3 及既有训练推理系列产品的全产品形态并兼容 TensorFlow 的 RealDiv 语义。本文以 math/real_div/README.md 为骨架结合算子 IR 定义、Shape 推导、AICore/AICPU 双通道 Kernel、L0 API 与单元测试源码完整讲解该算子的功能语义、参数约束、图模式调用方法以及底层实现原理帮助开发者在自定义算子开发、模型移植与性能分析场景中直接复用。功能说明逐元素实数除法与广播语义RealDiv 算子按元素逐个计算两个输入张量的商功能上等价于数值计算中的div计算公式为$$ y x_1 / x_2 $$其中x1为被除数张量x2为除数张量y为输出张量三者形状与数据类型保持一致布尔类型除外见下文参数说明。算子支持 NumPy 风格的广播broadcast运算当两个输入形状不完全一致时按从右向左对齐、维度为 1 的轴进行扩展最终输出形状为二者广播后的结果。从算子 IR 的官方注释看RealDiv 明确声明Returns x1/x2 element-wise for real types. Support broadcasting operations并且兼容 TensorFlow 的RealDiv算子见 real_div_proto.h。在实现层面广播不仅是图模式下的 Shape 推导能力更渗透到了 AICore 的广播调度模板与 AICPU 的 stride 广播迭代器中是理解本算子性能的关键。产品支持情况RealDiv 在 CANN ops-math 中面向以下产品形态提供支持产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品√Atlas 推理系列产品√Atlas 训练系列产品√上述产品清单与算子实现侧的配置一一对应算子注册时分别针对ascend950与ascend350添加了 AICore 配置见 real_div_def.cpp且op_host/config/目录下按平台分别维护了 ascend950 的算子配置 与 ascend350 的算子配置。在 AICore 路径的入口realdiv.cpp还会依据芯片版本如 ASCEND910B/910_93、ASCEND610LITE 及新平台动态选择数据类型支持清单具体差异在底层实现原理一节详述。参数说明RealDiv 算子共 3 个参数全部为张量无属性Attr参数参数名输入/输出/属性描述数据类型数据格式x1输入公式中的输入张量 x1被除数FLOAT, FLOAT16, BF16, DOUBLE, UINT8, INT8, UINT16, INT16, INT32, INT64, BOOL, COMPLEX64, COMPLEX128NDx2输入公式中的输入张量 x2除数同 x1NDy输出公式中的输出张量 y同 x1ND需要重点说明两点细节布尔输入的输出类型提升当x1为 BOOL 类型时输出y不再是 BOOL而是浮点类型。算子 IR 注释明确写道Has the same dtype and format as input x1 if the type of x1 is not bool. If the type of x1 is bool, y is float type.见 real_div_proto.h。这与 API 层实现一致在 realdiv.cpp 中当x1为DT_BOOL时输出张量被分配为DT_FLOAT。INT32 输入在特定平台上的输出提升对于 Ascend 910B/910_93 平台当两个输入均为 INT32 时输出同样会被提升为DT_FLOAT对应is910BInt32Supported分支以保证整数除法结果的精度表达。调用方式基于算子 IR 的图模式调用README 给出了 RealDiv 的标准调用路径——图模式GEIR调用通过 算子 IR 构图然后调用完整示例程序 test_geir_real_div.cpp。该示例完整展示了定义单算子图 → 初始化 GE 环境 → 构图与绑定输入 → 建图入会话 → 运行图 → 落盘输入输出 → 打印结果的完整流程可直接作为在 CANN 上以图模式运行任意单算子的最小模板。算子注册与构图入口示例通过op::RealDiv(realdivop)在图中创建算子节点并通过宏ADD_INPUT为x1、x2添加两个形状为{4, 4}的占位输入用ADD_OUTPUT声明同形状的输出y见 test_geir_real_div.cppint CreateOppInGraph(DataType inDtype, std::vectorge::Tensor input, std::vectorOperator inputs, std::vectorOperator outputs, Graph graph) { Status ret SUCCESS; // 自定义代码添加单算子定义到图中 auto realdivop op::RealDiv(realdivop); std::vectorint64_t xShape {4, 4}; ADD_INPUT(1, x1, inDtype, xShape); ADD_INPUT(2, x2, inDtype, xShape); ADD_OUTPUT(1, y, inDtype, xShape); outputs.push_back(realdivop); return SUCCESS; }其中ADD_INPUT宏使用op::Data(placeholderN)创建数据节点构造FORMAT_ND的TensorDesc将占位符所在节点set_input_x1/x2挂到 RealDiv 节点上并把占位节点同时加入输入列表ADD_OUTPUT宏则通过update_output_desc_y声明输出描述。示例默认以DT_FLOAT作为输入类型两个输入的数据均用GenOnesDataFloat32填充为常量 2.0。GE 初始化、构图与运行主函数依次完成以下步骤test_geir_real_div.cpp初始化 GE通过ge::GEInitialize(global_options)初始化图引擎global_options中设置ge.exec.deviceId 0与ge.graphRunMode 1。构图调用CreateOppInGraph生成单算子计算图随后graph.SetInputs(inputs).SetOutputs(outputs)将占位输入与算子输出挂为图的输入输出。创建会话并入图new Session(build_options)创建 GE 会话session-AddGraph(graph_id, graph, graph_options)将计算图加入会话。可选 dumpaclgrphDumpGraph(graph, ...)将构图结果导出到./dump目录便于检查 IR。运行图session-RunGraph(graph_id, input, output)执行计算示例中两个输入均为 2.0 的{4, 4}张量因此输出为全 1.0 的{4, 4}张量。落盘与打印输入输出分别写为tc_ge_irrun_test_0008_npu_input_{i}.bin与tc_ge_irrun_test_0008_npu_output_{i}.bin并把输出逐元素打印result[j]。收尾通过GEGetErrorMsgV2()/GEGetWarningMsgV2()获取执行告警与错误信息最后GEFinalize()释放环境。编译运行该示例需要 CANN 的 GEGraph Engine头文件graph.h、ge_api.h、array_ops.h、nn_other.h等与对应链接库并在 NPU 环境上执行。底层实现原理从算子定义到 AICore/AICPU 双通道RealDiv 的完整实现横跨五个层次算子 IRop_graph、算子定义与 Shape 推导op_host、AICore Kernelop_kernel/arch35、AICPU Kernelop_kernel_aicpu以及 L0 APIop_api。下面按调用链自底向上展开。算子 IR 定义REG_OP 注册与第三方框架兼容real_div_proto.h 使用 GE 的REG_OP宏注册 RealDiv 算子输入x1、x2与输出y的TensorType均覆盖 FLOAT、FLOAT16、BF16、DOUBLE、UINT8、INT8、UINT16、INT16、INT32、INT64、BOOL、COMPLEX64、COMPLEX128且输出y的类型集合不含 BOOL——这从 IR 层面落实了布尔输入输出提升为浮点的语义。文件头部注释明确标注该算子兼容 TensorFlow 的RealDiv这为 TF 模型迁移到 NPU 时的算子映射提供了依据。算子宿主侧定义平台配置、动态 shape 与精度收敛real_div_def.cpp 通过OpDef类完成算子宿主侧注册核心信息包括输入输出约束x1、x2、y均为必选参数数据格式限定为FORMAT_ND与 README 参数表一致同时给出UnknownShapeFormat以支持动态 shape 场景。AICore 配置aicoreConfig设置了DynamicCompileStaticFlag(true)动态编译静态化、DynamicRankSupportFlag(true)动态 rank 支持、DynamicShapeSupportFlag(true)动态 shape 支持、PrecisionReduceFlag(true)允许精度收敛等开关并指向real_div_apt的 kernel 文件该配置同时挂载到ascend950与ascend350两个平台。Shape 推导广播规则复用real_div_infershape.cpp 的推导实现极为简洁——直接调用框架层工具Ops::Base::InferShape4Broadcast(context)完成广播后输出形状的计算并在算子实现注册表中通过IMPL_OP_INFERSHAPE(RealDiv)挂接。这意味着 RealDiv 的广播语义与 CANN 框架内部统一的广播推导规则保持一致任何符合该规则的输入组合包括一个输入为标量、两输入 rank 不同等都能得到正确输出形状。对应地test_real_div_infershape.cpp 提供了宿主侧推导的单元测试。AICore Kernel按数据类型分派的高精度 DAG 模板AICore 侧入口 real_div_apt.cpp 是一个模板 kernel根据编译期确定的DTYPE_X1在__global__ __aicore__函数中经if constexpr分派到不同的 DAG 实现均复用Ops::Base::BroadcastSch广播调度模板先做 shape 感知的 tiling 切分再按 DAG 逐 tile 计算输入类型计算策略说明bool先 Cast 到 half 再 Cast 到 floatfloat 相除输出 float对应RealDivWithBoolbfloat16 / float16先 Cast 到 floatfloat 高精度相除结果 Cast 回原类型对应RealDivFloatWithCastfloat直接DivHighPrecision相除对应RealDivFloatWithoutCastint32输出 int32/ int64直接Vec::Div整数相除对应RealDivIntegerWithoutCastint32输出 float先 Cast 到 floatfloat 高精度相除输出 float对应RealDivIntegerWithCast这些 DAG 定义集中在 real_div_dag.h以类型别名组合CopyInBrc带广播的搬入、Cast类型转换、Div/DivHighPrecision除法、CopyOut搬出等原子算子。以 float16 为例完整数据流为x1、x2 广播搬入 → Cast 到 float → DivHighPrecision → Cast 回 half → CopyOut。这种低精度输入先升精度计算再收敛的策略正是PrecisionReduceFlag的落地体现保证了半精度/整型输入下的除法精度。tiling key 结构在 real_div_struct.h 中通过BRC_TEMP_SCH_MODE_KEY_DECL(schMode)声明宿主侧对应的 tiling 逻辑由 real_div_tiling_arch35.h 与 real_div_tiling_arch35.cpp 实现负责核数、UB 空间、workspace 的计算。AICPU Kernel广播场景下的 stride 优化迭代器当平台或数据类型不支持 AICore 路径时如 DOUBLE、COMPLEX64/128 等RealDiv 回退到 AICPU 执行实现在 real_div_aicpu.cpp。该 Kernel 按输入同 dtype前提分派RealDivSameTypeCompute覆盖从 FP16/FLOAT/DOUBLE、各整数类型到 COMPLEX64/128 的完整类型表同时注册了RealDiv与Div两个算子名OPS_MATH_REGISTER_CPU_KERNELV2。其性能核心是自定义的BcastDivInfostride 广播迭代器算法分四步补齐 rank 并校验广播合法性较短 shape 前补 1逐维校验相等或其一为 1否则返回不兼容计算有效 stride维度与输出一致时取自然 stride广播维度 stride 置 0使遍历时天然复用同一元素剔除 size 为 1 的维度无需迭代合并相邻连续维度当相邻两维在两个输入上同时满足连续性条件时折叠为单维降低循环深度。计算路径上对同 shapex1 单元素x2 单元素三种特殊情况走SpecialComputeImpl的专用循环一般广播场景则按最内层 stride 形态分支x_inner1 y_inner1时两输入最内层连续编译器可在 aarch64 上自动向量化为 NEONvdivqy_inner0y 最内层广播时把标量除数提升到循环外hoist避免重复访存其余情况走通用 strided 循环。此外片间并行通过CpuKernelUtils::ParallelFor实现按数据量动态选择核数与分片超过 16K/35K 元素阈值时提高并行度每个分片仅做一次 O(ndims) 的起始坐标分解随后以摊销 O(1) 的进位传播维护多维坐标——这是对朴素逐元素计算广播下标每元素 O(ndims) 除法的显著优化。AICPU 侧还包含除零保护对整数类型verify_zerotrue在计算前扫描除数张量发现 0 则报division by zero参数错误。L0 API 层平台感知的 AICore/AICPU 自动分派L0 API 头文件 realdiv.h 暴露了 4 组重载与 1 个查询函数realdiv.cpp 给出了完整实现RealDiv(self, other, executor)基础两输入除法先对两个ViewShape做BroadcastInferShape求出广播后输出形状DT_BOOL或 910B/910_93 的 INT32 输入输出提升为DT_FLOAT否则输出 dtype 同x1RealDivInplace(self, divisor, executor)原地版本要求self与divisor形状、dtype 完全一致输出直接写入divisorconst_cast后作为输出典型用于 torch 的div_语义RealDiv(self, other, mode, executor)mode为 0 时是 RealDivBOOL/INT32 输出提升为 FLOAT为 1 时是 TruncDiv 模式仅 BOOL 输出提升为 FLOATRealDiv(self, other, isScalar, executor)标量感知版本当isScalartrue且非 BOOL 时输出保持x1dtype否则提升为 FLOATIsRealDivSupportNonContiguous(self, other, outputShape)查询算子是否支持非连续输入用于上层框架决定是否先行contiguous()处理。分派逻辑的核心是IsAiCoreSupport它根据当前平台GetCurrentPlatformInfo().GetSocVersion()与输入 dtype 决定走 AICore 还是 AICPU。不同平台支持走 AICore 的类型清单存在差异例如默认平台为 FLOAT/FLOAT16/BF16/BOOL910B/910_93 额外支持 INT32610LITE 支持 FLOAT/FLOAT16/INT8/UINT8/INT32950 新平台则支持 FLOAT/FLOAT16/BF16/BOOL/INT32/INT64见 realdiv.cpp。若判定支持 AICore则通过ADD_TO_LAUNCHER_LIST_AICORE下发 AICore 任务否则经ADD_TO_LAUNCHER_LIST_AICPU走 AICPU 通道。测试与验证RealDiv 在 tests 下按层次组织测试L0 API 层test_aclnn_real_div.cpp覆盖 float、float16、bf16、int32、bool 五种 dtype以及不同 shape、1 维张量、标量广播{2,3}与{1}、空张量{0,3}、mode0/1、isScalartrue/false、IsRealDivSupportNonContiguous等 13 个用例每个用例均断言输出形状符合预期宿主侧test_real_div_infershape.cpp验证广播 Shape 推导AICPU 侧test_real_div.cpp直接驱动 CPU Kernel 验证计算正确性st 测试ttk_kernel_real_div_st.csv提供 AICore 通道的整算子基准用例配合 golden.py 生成预期结果。关键文件索引层次文件作用文档README.md功能、产品支持、参数与调用说明算子 IRreal_div_proto.hREG_OP定义与 TF RealDiv 兼容声明宿主定义real_div_def.cppOpDef 注册、AICore 平台配置Shape 推导real_div_infershape.cpp广播推导入口Tilingreal_div_tiling_arch35.harch35 tiling 结构AICore Kernelreal_div_apt.cpp / real_div_dag.h按 dtype 分派的 DAG 计算模板AICPU Kernelreal_div_aicpu.cpp广播优化的 CPU 实现L0 APIrealdiv.h / realdiv.cpp多形态入口与平台分派图模式示例test_geir_real_div.cpp完整可运行的 GE 构图调用示例单元测试test_aclnn_real_div.cpp 等多 dtype、广播与边界用例综上RealDiv 是一个在算子定义、Shape 推导、Kernel 实现与 L0 API 各层均已完备的成熟算子图模式调用路径可作为单算子开发的最小可运行模板AICore 侧低精度升精度计算与 AICPU 侧stride 广播 维度合并 并行分片则分别代表了 NPU 与 CPU 两种执行形态下的典型性能优化手法可作为后续自研除法类算子的直接参照。赞分享算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载相关推荐CANN ops-math 中 Atanh 算子的原理与调用指南从 aclnn 两段式接口到 AICore/AICPU 双实现CANN ops math 中 Atanh 算子的原理与调用指南从 aclnn 两段式接口到 AICore/AICPU 双实现 导读 Atanh反双曲正切算子库人工智能CANNCANN ops-math InvertPermutation 算子深度解析逆排列计算的图模式调用与 AICPU 实现CANN ops math InvertPermutation 算子深度解析逆排列计算的图模式调用与 AICPU 实现 本文围绕 CANN ops math算子库人工智能CANNCANN ops-math Eltwise 算子深度解析三种逐元素计算模式、参数配置与 GE 图模式调用实战CANN ops math Eltwise 算子深度解析三种逐元素计算模式、参数配置与 GE 图模式调用实战 Eltwise 是 CANN ops math算子库人工智能CANN上一篇5分钟快速掌握Mermaid Live Editor在线图表编辑终极指南下一篇Ralph for Claude Code 实战指南让自主开发循环把大项目拆成小任务、干完就停创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN ops-math RealDiv 算子深度解析：逐元素实数除法从图模式调用到 AICore/AICPU 双通道实现
算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载RealDiv 是 CANN ops-math 数学算子库中实现逐元素实数除法y x1 / x2的基础算子覆盖从 Ascend 950 到 Atlas A2/A3 及既有训练推理系列产品的全产品形态并兼容 TensorFlow 的 RealDiv 语义。本文以 math/real_div/README.md 为骨架结合算子 IR 定义、Shape 推导、AICore/AICPU 双通道 Kernel、L0 API 与单元测试源码完整讲解该算子的功能语义、参数约束、图模式调用方法以及底层实现原理帮助开发者在自定义算子开发、模型移植与性能分析场景中直接复用。功能说明逐元素实数除法与广播语义RealDiv 算子按元素逐个计算两个输入张量的商功能上等价于数值计算中的div计算公式为$$ y x_1 / x_2 $$其中x1为被除数张量x2为除数张量y为输出张量三者形状与数据类型保持一致布尔类型除外见下文参数说明。算子支持 NumPy 风格的广播broadcast运算当两个输入形状不完全一致时按从右向左对齐、维度为 1 的轴进行扩展最终输出形状为二者广播后的结果。从算子 IR 的官方注释看RealDiv 明确声明Returns x1/x2 element-wise for real types. Support broadcasting operations并且兼容 TensorFlow 的RealDiv算子见 real_div_proto.h。在实现层面广播不仅是图模式下的 Shape 推导能力更渗透到了 AICore 的广播调度模板与 AICPU 的 stride 广播迭代器中是理解本算子性能的关键。产品支持情况RealDiv 在 CANN ops-math 中面向以下产品形态提供支持产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品√Atlas 推理系列产品√Atlas 训练系列产品√上述产品清单与算子实现侧的配置一一对应算子注册时分别针对ascend950与ascend350添加了 AICore 配置见 real_div_def.cpp且op_host/config/目录下按平台分别维护了 ascend950 的算子配置 与 ascend350 的算子配置。在 AICore 路径的入口realdiv.cpp还会依据芯片版本如 ASCEND910B/910_93、ASCEND610LITE 及新平台动态选择数据类型支持清单具体差异在底层实现原理一节详述。参数说明RealDiv 算子共 3 个参数全部为张量无属性Attr参数参数名输入/输出/属性描述数据类型数据格式x1输入公式中的输入张量 x1被除数FLOAT, FLOAT16, BF16, DOUBLE, UINT8, INT8, UINT16, INT16, INT32, INT64, BOOL, COMPLEX64, COMPLEX128NDx2输入公式中的输入张量 x2除数同 x1NDy输出公式中的输出张量 y同 x1ND需要重点说明两点细节布尔输入的输出类型提升当x1为 BOOL 类型时输出y不再是 BOOL而是浮点类型。算子 IR 注释明确写道Has the same dtype and format as input x1 if the type of x1 is not bool. If the type of x1 is bool, y is float type.见 real_div_proto.h。这与 API 层实现一致在 realdiv.cpp 中当x1为DT_BOOL时输出张量被分配为DT_FLOAT。INT32 输入在特定平台上的输出提升对于 Ascend 910B/910_93 平台当两个输入均为 INT32 时输出同样会被提升为DT_FLOAT对应is910BInt32Supported分支以保证整数除法结果的精度表达。调用方式基于算子 IR 的图模式调用README 给出了 RealDiv 的标准调用路径——图模式GEIR调用通过 算子 IR 构图然后调用完整示例程序 test_geir_real_div.cpp。该示例完整展示了定义单算子图 → 初始化 GE 环境 → 构图与绑定输入 → 建图入会话 → 运行图 → 落盘输入输出 → 打印结果的完整流程可直接作为在 CANN 上以图模式运行任意单算子的最小模板。算子注册与构图入口示例通过op::RealDiv(realdivop)在图中创建算子节点并通过宏ADD_INPUT为x1、x2添加两个形状为{4, 4}的占位输入用ADD_OUTPUT声明同形状的输出y见 test_geir_real_div.cppint CreateOppInGraph(DataType inDtype, std::vectorge::Tensor input, std::vectorOperator inputs, std::vectorOperator outputs, Graph graph) { Status ret SUCCESS; // 自定义代码添加单算子定义到图中 auto realdivop op::RealDiv(realdivop); std::vectorint64_t xShape {4, 4}; ADD_INPUT(1, x1, inDtype, xShape); ADD_INPUT(2, x2, inDtype, xShape); ADD_OUTPUT(1, y, inDtype, xShape); outputs.push_back(realdivop); return SUCCESS; }其中ADD_INPUT宏使用op::Data(placeholderN)创建数据节点构造FORMAT_ND的TensorDesc将占位符所在节点set_input_x1/x2挂到 RealDiv 节点上并把占位节点同时加入输入列表ADD_OUTPUT宏则通过update_output_desc_y声明输出描述。示例默认以DT_FLOAT作为输入类型两个输入的数据均用GenOnesDataFloat32填充为常量 2.0。GE 初始化、构图与运行主函数依次完成以下步骤test_geir_real_div.cpp初始化 GE通过ge::GEInitialize(global_options)初始化图引擎global_options中设置ge.exec.deviceId 0与ge.graphRunMode 1。构图调用CreateOppInGraph生成单算子计算图随后graph.SetInputs(inputs).SetOutputs(outputs)将占位输入与算子输出挂为图的输入输出。创建会话并入图new Session(build_options)创建 GE 会话session-AddGraph(graph_id, graph, graph_options)将计算图加入会话。可选 dumpaclgrphDumpGraph(graph, ...)将构图结果导出到./dump目录便于检查 IR。运行图session-RunGraph(graph_id, input, output)执行计算示例中两个输入均为 2.0 的{4, 4}张量因此输出为全 1.0 的{4, 4}张量。落盘与打印输入输出分别写为tc_ge_irrun_test_0008_npu_input_{i}.bin与tc_ge_irrun_test_0008_npu_output_{i}.bin并把输出逐元素打印result[j]。收尾通过GEGetErrorMsgV2()/GEGetWarningMsgV2()获取执行告警与错误信息最后GEFinalize()释放环境。编译运行该示例需要 CANN 的 GEGraph Engine头文件graph.h、ge_api.h、array_ops.h、nn_other.h等与对应链接库并在 NPU 环境上执行。底层实现原理从算子定义到 AICore/AICPU 双通道RealDiv 的完整实现横跨五个层次算子 IRop_graph、算子定义与 Shape 推导op_host、AICore Kernelop_kernel/arch35、AICPU Kernelop_kernel_aicpu以及 L0 APIop_api。下面按调用链自底向上展开。算子 IR 定义REG_OP 注册与第三方框架兼容real_div_proto.h 使用 GE 的REG_OP宏注册 RealDiv 算子输入x1、x2与输出y的TensorType均覆盖 FLOAT、FLOAT16、BF16、DOUBLE、UINT8、INT8、UINT16、INT16、INT32、INT64、BOOL、COMPLEX64、COMPLEX128且输出y的类型集合不含 BOOL——这从 IR 层面落实了布尔输入输出提升为浮点的语义。文件头部注释明确标注该算子兼容 TensorFlow 的RealDiv这为 TF 模型迁移到 NPU 时的算子映射提供了依据。算子宿主侧定义平台配置、动态 shape 与精度收敛real_div_def.cpp 通过OpDef类完成算子宿主侧注册核心信息包括输入输出约束x1、x2、y均为必选参数数据格式限定为FORMAT_ND与 README 参数表一致同时给出UnknownShapeFormat以支持动态 shape 场景。AICore 配置aicoreConfig设置了DynamicCompileStaticFlag(true)动态编译静态化、DynamicRankSupportFlag(true)动态 rank 支持、DynamicShapeSupportFlag(true)动态 shape 支持、PrecisionReduceFlag(true)允许精度收敛等开关并指向real_div_apt的 kernel 文件该配置同时挂载到ascend950与ascend350两个平台。Shape 推导广播规则复用real_div_infershape.cpp 的推导实现极为简洁——直接调用框架层工具Ops::Base::InferShape4Broadcast(context)完成广播后输出形状的计算并在算子实现注册表中通过IMPL_OP_INFERSHAPE(RealDiv)挂接。这意味着 RealDiv 的广播语义与 CANN 框架内部统一的广播推导规则保持一致任何符合该规则的输入组合包括一个输入为标量、两输入 rank 不同等都能得到正确输出形状。对应地test_real_div_infershape.cpp 提供了宿主侧推导的单元测试。AICore Kernel按数据类型分派的高精度 DAG 模板AICore 侧入口 real_div_apt.cpp 是一个模板 kernel根据编译期确定的DTYPE_X1在__global__ __aicore__函数中经if constexpr分派到不同的 DAG 实现均复用Ops::Base::BroadcastSch广播调度模板先做 shape 感知的 tiling 切分再按 DAG 逐 tile 计算输入类型计算策略说明bool先 Cast 到 half 再 Cast 到 floatfloat 相除输出 float对应RealDivWithBoolbfloat16 / float16先 Cast 到 floatfloat 高精度相除结果 Cast 回原类型对应RealDivFloatWithCastfloat直接DivHighPrecision相除对应RealDivFloatWithoutCastint32输出 int32/ int64直接Vec::Div整数相除对应RealDivIntegerWithoutCastint32输出 float先 Cast 到 floatfloat 高精度相除输出 float对应RealDivIntegerWithCast这些 DAG 定义集中在 real_div_dag.h以类型别名组合CopyInBrc带广播的搬入、Cast类型转换、Div/DivHighPrecision除法、CopyOut搬出等原子算子。以 float16 为例完整数据流为x1、x2 广播搬入 → Cast 到 float → DivHighPrecision → Cast 回 half → CopyOut。这种低精度输入先升精度计算再收敛的策略正是PrecisionReduceFlag的落地体现保证了半精度/整型输入下的除法精度。tiling key 结构在 real_div_struct.h 中通过BRC_TEMP_SCH_MODE_KEY_DECL(schMode)声明宿主侧对应的 tiling 逻辑由 real_div_tiling_arch35.h 与 real_div_tiling_arch35.cpp 实现负责核数、UB 空间、workspace 的计算。AICPU Kernel广播场景下的 stride 优化迭代器当平台或数据类型不支持 AICore 路径时如 DOUBLE、COMPLEX64/128 等RealDiv 回退到 AICPU 执行实现在 real_div_aicpu.cpp。该 Kernel 按输入同 dtype前提分派RealDivSameTypeCompute覆盖从 FP16/FLOAT/DOUBLE、各整数类型到 COMPLEX64/128 的完整类型表同时注册了RealDiv与Div两个算子名OPS_MATH_REGISTER_CPU_KERNELV2。其性能核心是自定义的BcastDivInfostride 广播迭代器算法分四步补齐 rank 并校验广播合法性较短 shape 前补 1逐维校验相等或其一为 1否则返回不兼容计算有效 stride维度与输出一致时取自然 stride广播维度 stride 置 0使遍历时天然复用同一元素剔除 size 为 1 的维度无需迭代合并相邻连续维度当相邻两维在两个输入上同时满足连续性条件时折叠为单维降低循环深度。计算路径上对同 shapex1 单元素x2 单元素三种特殊情况走SpecialComputeImpl的专用循环一般广播场景则按最内层 stride 形态分支x_inner1 y_inner1时两输入最内层连续编译器可在 aarch64 上自动向量化为 NEONvdivqy_inner0y 最内层广播时把标量除数提升到循环外hoist避免重复访存其余情况走通用 strided 循环。此外片间并行通过CpuKernelUtils::ParallelFor实现按数据量动态选择核数与分片超过 16K/35K 元素阈值时提高并行度每个分片仅做一次 O(ndims) 的起始坐标分解随后以摊销 O(1) 的进位传播维护多维坐标——这是对朴素逐元素计算广播下标每元素 O(ndims) 除法的显著优化。AICPU 侧还包含除零保护对整数类型verify_zerotrue在计算前扫描除数张量发现 0 则报division by zero参数错误。L0 API 层平台感知的 AICore/AICPU 自动分派L0 API 头文件 realdiv.h 暴露了 4 组重载与 1 个查询函数realdiv.cpp 给出了完整实现RealDiv(self, other, executor)基础两输入除法先对两个ViewShape做BroadcastInferShape求出广播后输出形状DT_BOOL或 910B/910_93 的 INT32 输入输出提升为DT_FLOAT否则输出 dtype 同x1RealDivInplace(self, divisor, executor)原地版本要求self与divisor形状、dtype 完全一致输出直接写入divisorconst_cast后作为输出典型用于 torch 的div_语义RealDiv(self, other, mode, executor)mode为 0 时是 RealDivBOOL/INT32 输出提升为 FLOAT为 1 时是 TruncDiv 模式仅 BOOL 输出提升为 FLOATRealDiv(self, other, isScalar, executor)标量感知版本当isScalartrue且非 BOOL 时输出保持x1dtype否则提升为 FLOATIsRealDivSupportNonContiguous(self, other, outputShape)查询算子是否支持非连续输入用于上层框架决定是否先行contiguous()处理。分派逻辑的核心是IsAiCoreSupport它根据当前平台GetCurrentPlatformInfo().GetSocVersion()与输入 dtype 决定走 AICore 还是 AICPU。不同平台支持走 AICore 的类型清单存在差异例如默认平台为 FLOAT/FLOAT16/BF16/BOOL910B/910_93 额外支持 INT32610LITE 支持 FLOAT/FLOAT16/INT8/UINT8/INT32950 新平台则支持 FLOAT/FLOAT16/BF16/BOOL/INT32/INT64见 realdiv.cpp。若判定支持 AICore则通过ADD_TO_LAUNCHER_LIST_AICORE下发 AICore 任务否则经ADD_TO_LAUNCHER_LIST_AICPU走 AICPU 通道。测试与验证RealDiv 在 tests 下按层次组织测试L0 API 层test_aclnn_real_div.cpp覆盖 float、float16、bf16、int32、bool 五种 dtype以及不同 shape、1 维张量、标量广播{2,3}与{1}、空张量{0,3}、mode0/1、isScalartrue/false、IsRealDivSupportNonContiguous等 13 个用例每个用例均断言输出形状符合预期宿主侧test_real_div_infershape.cpp验证广播 Shape 推导AICPU 侧test_real_div.cpp直接驱动 CPU Kernel 验证计算正确性st 测试ttk_kernel_real_div_st.csv提供 AICore 通道的整算子基准用例配合 golden.py 生成预期结果。关键文件索引层次文件作用文档README.md功能、产品支持、参数与调用说明算子 IRreal_div_proto.hREG_OP定义与 TF RealDiv 兼容声明宿主定义real_div_def.cppOpDef 注册、AICore 平台配置Shape 推导real_div_infershape.cpp广播推导入口Tilingreal_div_tiling_arch35.harch35 tiling 结构AICore Kernelreal_div_apt.cpp / real_div_dag.h按 dtype 分派的 DAG 计算模板AICPU Kernelreal_div_aicpu.cpp广播优化的 CPU 实现L0 APIrealdiv.h / realdiv.cpp多形态入口与平台分派图模式示例test_geir_real_div.cpp完整可运行的 GE 构图调用示例单元测试test_aclnn_real_div.cpp 等多 dtype、广播与边界用例综上RealDiv 是一个在算子定义、Shape 推导、Kernel 实现与 L0 API 各层均已完备的成熟算子图模式调用路径可作为单算子开发的最小可运行模板AICore 侧低精度升精度计算与 AICPU 侧stride 广播 维度合并 并行分片则分别代表了 NPU 与 CPU 两种执行形态下的典型性能优化手法可作为后续自研除法类算子的直接参照。赞分享算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载相关推荐CANN ops-math 中 Atanh 算子的原理与调用指南从 aclnn 两段式接口到 AICore/AICPU 双实现CANN ops math 中 Atanh 算子的原理与调用指南从 aclnn 两段式接口到 AICore/AICPU 双实现 导读 Atanh反双曲正切算子库人工智能CANNCANN ops-math InvertPermutation 算子深度解析逆排列计算的图模式调用与 AICPU 实现CANN ops math InvertPermutation 算子深度解析逆排列计算的图模式调用与 AICPU 实现 本文围绕 CANN ops math算子库人工智能CANNCANN ops-math Eltwise 算子深度解析三种逐元素计算模式、参数配置与 GE 图模式调用实战CANN ops math Eltwise 算子深度解析三种逐元素计算模式、参数配置与 GE 图模式调用实战 Eltwise 是 CANN ops math算子库人工智能CANN上一篇5分钟快速掌握Mermaid Live Editor在线图表编辑终极指南下一篇Ralph for Claude Code 实战指南让自主开发循环把大项目拆成小任务、干完就停创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考