尧图网页设计 
CANN ops-math Fills 算子详解NPU 上基于图模式实现 Tensor 全量填充【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math导读本文以 CANN ops-math 开源仓库中 conversion/fills/README.md 为核心系统讲解数学基础算子库中的Fills 算子它将输入 Tensor 的所有元素统一填充为指定值等价于 PyTorch 的Tensor.fill_。文章将覆盖产品支持范围、参数语义、算子原型与注册实现并结合仓库源码examples/test_geir_fills.cpp、op_host/fills_def.cpp、op_kernel/fills_apt.cpp等深入剖析 InferShape、Tiling、Kernel 的执行链路与图模式调用实战方法帮助你快速在 NPUAscend环境中部署和验证 Fills 算子。产品支持情况Fills 算子已在当前仓库中注册并适配多款昇腾硬件产品支持的矩阵如下产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品√Atlas 推理系列产品√Atlas 训练系列产品√从源码看该适配范围与算子注册配置一一对应op_host/fills_def.cpp 中通过this-AICore().AddConfig(ascend950, aicoreConfig)与this-AICore().AddConfig(ascend350, aicoreConfig)分别为 ascend950对应 Ascend 950 系列与 ascend350对应 Atlas A3 等采用 arch35 架构的产品添加了 AICore 配置同时在 op_host/config/ascend950/ 与 op_host/config/ascend350/ 下提供了对应的fills_binary.json与fills_simplified_key.ini注册清单。功能说明算子功能将输入 Tensorx的所有元素填充为指定值value得到输出 Tensory。计算公式$$ out_i value $$即输出张量的每一个元素都与填充值相等输出与输入的 Shape 和数据类型保持一致。从算子原型注释op_graph/fills_proto.h可知该算子与 PyTorch 的fills算子兼容属于逐元素elewise类基础算子Kernel 侧正是通过 elewise 调度框架ElementwiseSch实现的。参数说明Fills 算子共包含两个输入/输出参数和一个属性参数官方参数定义如下参数名输入/输出/属性描述数据类型数据格式x输入待进行 Fills 计算的入参。FLOAT、FLOAT16、BFLOAT16、INT8、UINT8、INT32、INT64NDvalue输入属性待进行 Fills 填充值公式中的 value。FLOATNDy输出待进行 Fills 计算的出参公式中的 out_i。FLOAT、FLOAT16、BFLOAT16、INT8、UINT8、INT32、INT64ND参数语义的源码级印证x/y数据类型与 op_host/fills_def.cpp 中 OpDef 注册的数据类型集合完全一致DT_FLOAT16, DT_BF16, DT_FLOAT, DT_INT8, DT_UINT8, DT_INT32, DT_INT64。需要说明的是算子原型fills_proto.h中REG_OP(Fills)的TensorType为{BasicType(), DT_BOOL}即底层 proto 层面还额外声明了 BOOL 类型兼容但当前 host 侧 OpDef 与 Kernel/Tiling 的实际支持以 README 列出的 7 种类型为准。value属性类型原型中声明为REQUIRED_ATTR(value, Float)即必选REQUIRED的 FLOAT 类型属性Tiling 阶段通过attrs-GetAttrPointerfloat(0)读取。数据格式输入、输出均为ND 格式注册中的Format与UnknownShapeFormat全部为FORMAT_ND说明 Fills 算子面向的是最通用的稠密 N 维张量布局。约束说明官方 README 中明确约束说明无。即 Fills 算子没有额外的使用限制条件。结合源码可以补充两点隐含约束均在 Tiling 阶段校验见 fills_tiling_arch35.cpp输入输出数据类型必须一致CalcOutputDtype()中校验inputDtype outputDtype否则报错 The dtypes of x(input) and y(output) must be the same。输入输出 Shape 必须一致CheckShape()中校验inputYShape outputZShape否则报错 The shapes of x(input) and y(output) must be the same输出张量逐元素填充Shape 自然与输入一致。图模式调用说明Fills 算子官方推荐的调用方式是图模式Graph Mode即通过算子 IRFills的 GE 算子原型构建计算图后提交执行。官方调用方式对照如下调用方式调用样例说明图模式调用conversion/fills/examples/test_geir_fills.cpp通过算子IR构图方式调用 Fills 算子。示例整体流程test_geir_fills.cpp是一个完整的单算子图执行样例主流程main分为五个阶段初始化 GE 环境InitEnv通过ge::GEInitialize设置全局选项例如{ge.exec.deviceId, 0}, {ge.graphRunMode, 1}创建计算图并配置输入输出CreateAndConfigGraph→CreateOppInGraph创建 Session 并将图加入会话AddGraphToSessionsession-AddGraph(graph_id, graph, ...)执行图DumpAndRunGraphsession-RunGraph(graph_id, input, output)并可调用aclgrphDumpGraph将图 dump 到本地./dump目录以便排查处理输入输出数据并清理资源ProcessInputData/ProcessOutputData/FinalizeRes。核心构图代码讲解在CreateOppInGraph中通过op::Fills创建算子节点并使用宏完成输入输出与属性的装配// 自定义代码添加单算子定义到图中 auto add1 op::Fills(fills); vectorvectorint64_t shape {3, 4}; ADD_INPUT(1, x, inDtype, shape); // 创建 shape 为 {3, 4} 的 ND 输入占位节点数据类型 inDtype ADD_OUTPUT(1, y, inDtype, shape); // 声明输出 y数据类型与 shape 与输入一致 ADD_INPUT_ATTR(value, 1.5f); // 设置必选属性 value 1.5f即填充值其中ADD_INPUT宏展开后做了四件事用op::Data(placeholder1)创建输入占位算子并构造TensorDesc(ge::Shape({3,4}), ge::FORMAT_ND, inDtype)通过GenOnesData按 Shape 生成全 2 的初始数据仅作为运行时输入的载体将占位算子加入图中graph.AddOp(...)并用add1.set_input_x(placeholder1)建立 Fills 的输入连接把生成的Tensor存入input向量作为后续RunGraph的实际输入数据。示例中在CreateAndConfigGraph里指定DataType inDtype DT_FLOAT即 FLOAT32随后if (!inputs.empty() !outputs.empty()) { graph.SetInputs(inputs).SetOutputs(outputs); }将图输入输出与 Session 关联。执行结束后ProcessOutputData会把结果写入./tc_ge_irrun_test_0008_npu_output_0.bin并逐元素打印result[j]便于直接核对输出是否全部等于value。运行与验证建议该示例依赖 GEGraph Engine运行环境与ge_ir_build、nn_other等头文件编译时需链接 CANN 的 GE 库并在装有昇腾设备如 Ascend 950 / Atlas A3的环境上执行由于填充结果是确定性的验证时只需检查输出文件中的所有元素是否等于属性value示例中为 1.5ST 测试亦采用类似的数值比对方式见下文测试与验证小节。算子实现的源码级剖析1. 算子原型IR 定义op_graph/fills_proto.h 使用REG_OP宏声明了算子的 IR 原型REG_OP(Fills) .INPUT(x, TensorType({BasicType(), DT_BOOL})) .OUTPUT(y, TensorType({BasicType(), DT_BOOL})) .REQUIRED_ATTR(value, Float) .OP_END_FACTORY_REG(Fills)同时该头文件还记录了关键的精度说明对于 FLOAT32 类型参数无精度损失对于 INT32 和 INT64 参数当参数值超过 2^24 时会发生精度损失此时建议改用 Fill 算子。这是选择value取值时需要特别关注的边界value属性本身是 FLOAT转成 INT32/INT64 输出时超出 2^24 的整数值无法被精确表达。2. 算子注册OpDefop_host/fills_def.cpp 通过OpDef类完成 host 侧注册除前面提到的输入输出类型/格式外还声明了 AICore 配置OpAICoreConfig aicoreConfig; aicoreConfig.DynamicCompileStaticFlag(true) .DynamicFormatFlag(false) .DynamicRankSupportFlag(true) // 支持动态 Rank .DynamicShapeSupportFlag(true) // 支持动态 Shape .NeedCheckSupportFlag(false) .PrecisionReduceFlag(true) .ExtendCfgInfo(opFile.value, fills_apt); // 指定 kernel 文件 fills_apt this-AICore().AddConfig(ascend950, aicoreConfig); this-AICore().AddConfig(ascend350, aicoreConfig);可以看出 Fills 是一个支持动态 Shape 与动态 Rank的算子且采用fills_apt作为 kernel 文件对应 op_kernel/fills_apt.cpp。3. 形状与数据类型推导InferShapeop_host/fills_infershape.cpp 直接复用逐元素算子的通用推导Ops::Base::InferShape4Elewise即输出 Shape 与输入一致。InferDataTypeop_graph/fills_graph_infer.cpp 中InferDataType4Fills将输出数据类型设置为与输入x相同SetOutputDataType(0, input_x_dtype)保证y与x类型一致。4. Tiling 策略op_host/arch35/fills_tiling_arch35.cpp 实现了 arch35ascend350/ascend950平台的 Tiling核心逻辑在RunTiling中CalcInputDtype/CalcOutputDtype校验数据类型并缓存inputDtype/outputDtypeCheckShape校验输入输出 Shape 一致SetAttr读取 FLOAT 属性value再根据outputDtype通过联合体FillsValue将其转换为对应类型如fp16_t、bfloat16、int8_t、int32_t、int64_t等写入tiling-valueElewiseBaseTiling.DoTilingFillsDAGT::OpDag按输出类型实例化 DAG 并生成基础分块baseTilingSetTilingData根据 dtype 映射tilingKey如FILLS_TPL_FP16、FILLS_TPL_FP32等设置 block 维度并申请约 16 MBASCEND_WORKSPACE 16777216的系统 workspace。Tiling 数据结构定义在 op_kernel/arch35/fills_tilingdata.hstruct FillsTilingData { EleBaseTilingData baseTiling; int64_t value; // 转换后的填充值以 8 字节形式承载各类型的位模式 };5. Kernel 执行op_kernel/fills_apt.cpp 是算子核函数实现基于 AscendC 编程模型template uint64_t schMode, uint64_t dType __global__ __aicore__ void fills(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); REGISTER_TILING_DEFAULT(FillsTilingData); GET_TILING_DATA(tilingData, tiling); TPipe pipe; ElementwiseSchschMode, FillsDAGDTYPE_X::OpDag sch((tilingData.baseTiling), pipe); sch.template SetVarDTYPE_X, 0(*(DTYPE_X*)(tilingData.value)); sch.Init(y); sch.Process(); }核心填充逻辑被封装成计算图 DAGop_kernel/arch35/fills_dag.htemplate typename T struct FillsDAG { using OpDuplicate BindVec::DuplicateT, Placeholder::VarT, 0; // 将 value 复制为向量 using OpCopyOut BindVec::CopyOutT, Placeholder::Out0T, OpDuplicate; // 写出到输出 using Outputs ElemsOpCopyOut; using OpDag DAGSchOutputs; };即先由Vec::Duplicate把标量填充值复制成向量再由Vec::CopyOut写出到输出张量最终由ElementwiseSch依据baseTiling的分块信息迭代整张输出。这也从实现层面印证了公式out_i valueFills 不读取输入数据内容仅按输入 Shape 将输出铺满 value。测试与验证仓库为 Fills 算子提供了完整的单测与系统测试用例ST 用例conversion/fills/tests/st/arch35/ttk_kernel_fills_st.csv 包含多个真实网络场景的用例如 EfficientNet-B0、FFDNet、TransformerXL、SimplePose、EfficientNet-B2覆盖 float16 / float32 输入Shape 从 1 维如(133,)到 5 维如(4, 4, 50, 128, 128)value取 0.0、3.90625e-3、5.9171598e-5 等典型值输入数据范围(-1, 1)float16 精度容差 1e-03、float32 精度容差 1e-04并设置绝对精度 1e-08。UT 用例conversion/fills/tests/ut/op_host/test_fills_infershape.cpp 与 conversion/fills/tests/ut/op_host/arch35/test_fills_tiling_arch35.cpp 分别覆盖 InferShape 与 Tiling 逻辑可在不依赖真实设备的情况下先行验证 shape 推导与 tiling 结果。总结Fills 算子是一个语义极简out_i value、实现路径清晰的数学基础算子上层通过 fills_proto.h 的 IR 原型与 fills_def.cpp 完成注册经 elewise 通用 InferShape 与显式 InferDataType 推导再由 arch35 Tiling 完成 dtype 转换与分块规划最终由ElementwiseSch FillsDAG在 NPU 上高效铺满输出。开发者既可以直接参照 test_geir_fills.cpp 以图模式在真实设备上运行验证也可以基于仓库内的 ST/UT 用例快速做回归测试需要留意 INT32/INT64 场景下value超过 2^24 的精度损失边界必要时按原型注释改用 Fill 算子。【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN ops-math Fills 算子详解：NPU 上基于图模式实现 Tensor 全量填充
CANN ops-math Fills 算子详解NPU 上基于图模式实现 Tensor 全量填充【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math导读本文以 CANN ops-math 开源仓库中 conversion/fills/README.md 为核心系统讲解数学基础算子库中的Fills 算子它将输入 Tensor 的所有元素统一填充为指定值等价于 PyTorch 的Tensor.fill_。文章将覆盖产品支持范围、参数语义、算子原型与注册实现并结合仓库源码examples/test_geir_fills.cpp、op_host/fills_def.cpp、op_kernel/fills_apt.cpp等深入剖析 InferShape、Tiling、Kernel 的执行链路与图模式调用实战方法帮助你快速在 NPUAscend环境中部署和验证 Fills 算子。产品支持情况Fills 算子已在当前仓库中注册并适配多款昇腾硬件产品支持的矩阵如下产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品√Atlas 推理系列产品√Atlas 训练系列产品√从源码看该适配范围与算子注册配置一一对应op_host/fills_def.cpp 中通过this-AICore().AddConfig(ascend950, aicoreConfig)与this-AICore().AddConfig(ascend350, aicoreConfig)分别为 ascend950对应 Ascend 950 系列与 ascend350对应 Atlas A3 等采用 arch35 架构的产品添加了 AICore 配置同时在 op_host/config/ascend950/ 与 op_host/config/ascend350/ 下提供了对应的fills_binary.json与fills_simplified_key.ini注册清单。功能说明算子功能将输入 Tensorx的所有元素填充为指定值value得到输出 Tensory。计算公式$$ out_i value $$即输出张量的每一个元素都与填充值相等输出与输入的 Shape 和数据类型保持一致。从算子原型注释op_graph/fills_proto.h可知该算子与 PyTorch 的fills算子兼容属于逐元素elewise类基础算子Kernel 侧正是通过 elewise 调度框架ElementwiseSch实现的。参数说明Fills 算子共包含两个输入/输出参数和一个属性参数官方参数定义如下参数名输入/输出/属性描述数据类型数据格式x输入待进行 Fills 计算的入参。FLOAT、FLOAT16、BFLOAT16、INT8、UINT8、INT32、INT64NDvalue输入属性待进行 Fills 填充值公式中的 value。FLOATNDy输出待进行 Fills 计算的出参公式中的 out_i。FLOAT、FLOAT16、BFLOAT16、INT8、UINT8、INT32、INT64ND参数语义的源码级印证x/y数据类型与 op_host/fills_def.cpp 中 OpDef 注册的数据类型集合完全一致DT_FLOAT16, DT_BF16, DT_FLOAT, DT_INT8, DT_UINT8, DT_INT32, DT_INT64。需要说明的是算子原型fills_proto.h中REG_OP(Fills)的TensorType为{BasicType(), DT_BOOL}即底层 proto 层面还额外声明了 BOOL 类型兼容但当前 host 侧 OpDef 与 Kernel/Tiling 的实际支持以 README 列出的 7 种类型为准。value属性类型原型中声明为REQUIRED_ATTR(value, Float)即必选REQUIRED的 FLOAT 类型属性Tiling 阶段通过attrs-GetAttrPointerfloat(0)读取。数据格式输入、输出均为ND 格式注册中的Format与UnknownShapeFormat全部为FORMAT_ND说明 Fills 算子面向的是最通用的稠密 N 维张量布局。约束说明官方 README 中明确约束说明无。即 Fills 算子没有额外的使用限制条件。结合源码可以补充两点隐含约束均在 Tiling 阶段校验见 fills_tiling_arch35.cpp输入输出数据类型必须一致CalcOutputDtype()中校验inputDtype outputDtype否则报错 The dtypes of x(input) and y(output) must be the same。输入输出 Shape 必须一致CheckShape()中校验inputYShape outputZShape否则报错 The shapes of x(input) and y(output) must be the same输出张量逐元素填充Shape 自然与输入一致。图模式调用说明Fills 算子官方推荐的调用方式是图模式Graph Mode即通过算子 IRFills的 GE 算子原型构建计算图后提交执行。官方调用方式对照如下调用方式调用样例说明图模式调用conversion/fills/examples/test_geir_fills.cpp通过算子IR构图方式调用 Fills 算子。示例整体流程test_geir_fills.cpp是一个完整的单算子图执行样例主流程main分为五个阶段初始化 GE 环境InitEnv通过ge::GEInitialize设置全局选项例如{ge.exec.deviceId, 0}, {ge.graphRunMode, 1}创建计算图并配置输入输出CreateAndConfigGraph→CreateOppInGraph创建 Session 并将图加入会话AddGraphToSessionsession-AddGraph(graph_id, graph, ...)执行图DumpAndRunGraphsession-RunGraph(graph_id, input, output)并可调用aclgrphDumpGraph将图 dump 到本地./dump目录以便排查处理输入输出数据并清理资源ProcessInputData/ProcessOutputData/FinalizeRes。核心构图代码讲解在CreateOppInGraph中通过op::Fills创建算子节点并使用宏完成输入输出与属性的装配// 自定义代码添加单算子定义到图中 auto add1 op::Fills(fills); vectorvectorint64_t shape {3, 4}; ADD_INPUT(1, x, inDtype, shape); // 创建 shape 为 {3, 4} 的 ND 输入占位节点数据类型 inDtype ADD_OUTPUT(1, y, inDtype, shape); // 声明输出 y数据类型与 shape 与输入一致 ADD_INPUT_ATTR(value, 1.5f); // 设置必选属性 value 1.5f即填充值其中ADD_INPUT宏展开后做了四件事用op::Data(placeholder1)创建输入占位算子并构造TensorDesc(ge::Shape({3,4}), ge::FORMAT_ND, inDtype)通过GenOnesData按 Shape 生成全 2 的初始数据仅作为运行时输入的载体将占位算子加入图中graph.AddOp(...)并用add1.set_input_x(placeholder1)建立 Fills 的输入连接把生成的Tensor存入input向量作为后续RunGraph的实际输入数据。示例中在CreateAndConfigGraph里指定DataType inDtype DT_FLOAT即 FLOAT32随后if (!inputs.empty() !outputs.empty()) { graph.SetInputs(inputs).SetOutputs(outputs); }将图输入输出与 Session 关联。执行结束后ProcessOutputData会把结果写入./tc_ge_irrun_test_0008_npu_output_0.bin并逐元素打印result[j]便于直接核对输出是否全部等于value。运行与验证建议该示例依赖 GEGraph Engine运行环境与ge_ir_build、nn_other等头文件编译时需链接 CANN 的 GE 库并在装有昇腾设备如 Ascend 950 / Atlas A3的环境上执行由于填充结果是确定性的验证时只需检查输出文件中的所有元素是否等于属性value示例中为 1.5ST 测试亦采用类似的数值比对方式见下文测试与验证小节。算子实现的源码级剖析1. 算子原型IR 定义op_graph/fills_proto.h 使用REG_OP宏声明了算子的 IR 原型REG_OP(Fills) .INPUT(x, TensorType({BasicType(), DT_BOOL})) .OUTPUT(y, TensorType({BasicType(), DT_BOOL})) .REQUIRED_ATTR(value, Float) .OP_END_FACTORY_REG(Fills)同时该头文件还记录了关键的精度说明对于 FLOAT32 类型参数无精度损失对于 INT32 和 INT64 参数当参数值超过 2^24 时会发生精度损失此时建议改用 Fill 算子。这是选择value取值时需要特别关注的边界value属性本身是 FLOAT转成 INT32/INT64 输出时超出 2^24 的整数值无法被精确表达。2. 算子注册OpDefop_host/fills_def.cpp 通过OpDef类完成 host 侧注册除前面提到的输入输出类型/格式外还声明了 AICore 配置OpAICoreConfig aicoreConfig; aicoreConfig.DynamicCompileStaticFlag(true) .DynamicFormatFlag(false) .DynamicRankSupportFlag(true) // 支持动态 Rank .DynamicShapeSupportFlag(true) // 支持动态 Shape .NeedCheckSupportFlag(false) .PrecisionReduceFlag(true) .ExtendCfgInfo(opFile.value, fills_apt); // 指定 kernel 文件 fills_apt this-AICore().AddConfig(ascend950, aicoreConfig); this-AICore().AddConfig(ascend350, aicoreConfig);可以看出 Fills 是一个支持动态 Shape 与动态 Rank的算子且采用fills_apt作为 kernel 文件对应 op_kernel/fills_apt.cpp。3. 形状与数据类型推导InferShapeop_host/fills_infershape.cpp 直接复用逐元素算子的通用推导Ops::Base::InferShape4Elewise即输出 Shape 与输入一致。InferDataTypeop_graph/fills_graph_infer.cpp 中InferDataType4Fills将输出数据类型设置为与输入x相同SetOutputDataType(0, input_x_dtype)保证y与x类型一致。4. Tiling 策略op_host/arch35/fills_tiling_arch35.cpp 实现了 arch35ascend350/ascend950平台的 Tiling核心逻辑在RunTiling中CalcInputDtype/CalcOutputDtype校验数据类型并缓存inputDtype/outputDtypeCheckShape校验输入输出 Shape 一致SetAttr读取 FLOAT 属性value再根据outputDtype通过联合体FillsValue将其转换为对应类型如fp16_t、bfloat16、int8_t、int32_t、int64_t等写入tiling-valueElewiseBaseTiling.DoTilingFillsDAGT::OpDag按输出类型实例化 DAG 并生成基础分块baseTilingSetTilingData根据 dtype 映射tilingKey如FILLS_TPL_FP16、FILLS_TPL_FP32等设置 block 维度并申请约 16 MBASCEND_WORKSPACE 16777216的系统 workspace。Tiling 数据结构定义在 op_kernel/arch35/fills_tilingdata.hstruct FillsTilingData { EleBaseTilingData baseTiling; int64_t value; // 转换后的填充值以 8 字节形式承载各类型的位模式 };5. Kernel 执行op_kernel/fills_apt.cpp 是算子核函数实现基于 AscendC 编程模型template uint64_t schMode, uint64_t dType __global__ __aicore__ void fills(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); REGISTER_TILING_DEFAULT(FillsTilingData); GET_TILING_DATA(tilingData, tiling); TPipe pipe; ElementwiseSchschMode, FillsDAGDTYPE_X::OpDag sch((tilingData.baseTiling), pipe); sch.template SetVarDTYPE_X, 0(*(DTYPE_X*)(tilingData.value)); sch.Init(y); sch.Process(); }核心填充逻辑被封装成计算图 DAGop_kernel/arch35/fills_dag.htemplate typename T struct FillsDAG { using OpDuplicate BindVec::DuplicateT, Placeholder::VarT, 0; // 将 value 复制为向量 using OpCopyOut BindVec::CopyOutT, Placeholder::Out0T, OpDuplicate; // 写出到输出 using Outputs ElemsOpCopyOut; using OpDag DAGSchOutputs; };即先由Vec::Duplicate把标量填充值复制成向量再由Vec::CopyOut写出到输出张量最终由ElementwiseSch依据baseTiling的分块信息迭代整张输出。这也从实现层面印证了公式out_i valueFills 不读取输入数据内容仅按输入 Shape 将输出铺满 value。测试与验证仓库为 Fills 算子提供了完整的单测与系统测试用例ST 用例conversion/fills/tests/st/arch35/ttk_kernel_fills_st.csv 包含多个真实网络场景的用例如 EfficientNet-B0、FFDNet、TransformerXL、SimplePose、EfficientNet-B2覆盖 float16 / float32 输入Shape 从 1 维如(133,)到 5 维如(4, 4, 50, 128, 128)value取 0.0、3.90625e-3、5.9171598e-5 等典型值输入数据范围(-1, 1)float16 精度容差 1e-03、float32 精度容差 1e-04并设置绝对精度 1e-08。UT 用例conversion/fills/tests/ut/op_host/test_fills_infershape.cpp 与 conversion/fills/tests/ut/op_host/arch35/test_fills_tiling_arch35.cpp 分别覆盖 InferShape 与 Tiling 逻辑可在不依赖真实设备的情况下先行验证 shape 推导与 tiling 结果。总结Fills 算子是一个语义极简out_i value、实现路径清晰的数学基础算子上层通过 fills_proto.h 的 IR 原型与 fills_def.cpp 完成注册经 elewise 通用 InferShape 与显式 InferDataType 推导再由 arch35 Tiling 完成 dtype 转换与分块规划最终由ElementwiseSch FillsDAG在 NPU 上高效铺满输出。开发者既可以直接参照 test_geir_fills.cpp 以图模式在真实设备上运行验证也可以基于仓库内的 ST/UT 用例快速做回归测试需要留意 INT32/INT64 场景下value超过 2^24 的精度损失边界必要时按原型注释改用 Fill 算子。【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考