算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载导读Invert 是 CANN ops-math 数学算子库中的按位取反算子对输入张量的每一个元素执行按位取反bitwise NOT运算等价于out_i ~input_i。本文以 math/invert/README.md 为核心骨架结合算子注册、构图、Tiling、Kernel 与单测等源码完整梳理 Invert 算子在 Ascend NPU 上的支持范围、参数语义、底层实现原理与图模式调用方法。读完本文你将掌握Invert 算子的数据类型/产品支持矩阵、逐元素算子的标准实现范式Proto → Def → Infershape → Tiling → Kernel、以及如何用算子 IR 构图方式在 GE 会话中跑通单算子图。一、算子功能与数学定义Invert 是典型的逐元素elementwise / elewise一元算子功能定义非常简洁算子功能对输入张量x的每一个元素执行按位取反bitwise invert得到输出张量y。计算公式$$ out_i\sim input_i $$即每个输出元素是输入元素按位取反后的结果。以 INT8 为例输入0b000000011经 invert 后输出0b11111110-2输入0b11110000则输出0b00001111。由于是按位运算invert(invert(x)) x恒成立因此该算子在符号/掩码翻转、编码反转等场景中常被用作基础原语。需要特别说明的是这里的取反是按位取反bitwise NOT不是算术取负negation-x。算术取负在 CANN 中由独立的 neg 算子承担见 math/neg/README.md两者语义不同~x在补码表示下等价于-x - 1。二、产品支持情况Invert 算子在各产品形态上的支持情况如下表取自 math/invert/README.md产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品√Atlas 训练系列产品√从源码可以交叉印证该支持矩阵。在 op_host/invert_def.cpp 中算子定义通过this-AICore().AddConfig(ascend950, aicoreConfig)显式注册了 ascend950 架构对应 Ascend 950 系列的 AICore 配置在 op_api/invert.cpp 中算子运行时会根据当前 NPU 架构动态选择执行路径// math/invert/op_api/invert.cpp static inline bool IsAiCoreSupport(const aclTensor* self) { auto npuArch op::GetCurrentPlatformInfo().GetCurNpuArch(); if (IsRegBase(npuArch)) { return CheckType(self-GetDataType(), REGBASE_DTYPE_SUPPORT_LIST); } return CheckType(self-GetDataType(), AICORE_DTYPE_SUPPORT_LIST); }即RegBase 类架构如 A2/A3 等通用寄存器基架构支持 8 种整数类型全部走 AICore其余架构仅 INT16、UINT16 两种类型走 AICore其余类型回退到 AICpu 执行。这也解释了为什么 README 中标注 Atlas 200I/500 A2 推理产品不支持——该形态无对应的 AICore/AICpu 内核路径。注意README 中不支持 BFLOAT16的结论在参数表中已明确体现Invert 仅接受 8 种整数类型浮点与 BF16 均不在支持范围内。三、参数说明Invert 算子只有一对输入输出参数取自 math/invert/README.md参数名输入/输出/属性描述数据类型数据格式x输入待进行 invert 计算的入参公式中的 input_iINT8、INT16、INT32、INT64、UINT8、UINT16、UINT32、UINT64NDy输出待进行 invert 计算的出参公式中的 out_iINT8、INT16、INT32、INT64、UINT8、UINT16、UINT32、UINT64ND附加约束Atlas 训练系列产品、Atlas 推理系列产品不支持 BFLOAT16。约束说明无除数据类型约束外Invert 不引入其他 shape、format 层面的限制。3.1 参数语义在算子定义中的落地参数声明与 README 完全对应见 op_host/invert_def.cppthis-Input(x) .ParamType(REQUIRED) .DataType({ge::DT_INT8, ge::DT_INT16, ge::DT_INT32, ge::DT_INT64, ge::DT_UINT8, ge::DT_UINT16, ge::DT_UINT32, ge::DT_UINT64}) .Format({ge::FORMAT_ND, ...}); this-Output(y) .ParamType(REQUIRED) .DataType({...}) // 与 x 完全一致 .Format({ge::FORMAT_ND, ...});该定义同时设置了 AICore 配置项DynamicCompileStaticFlag(true)支持动态编译静态化DynamicFormatFlag(false)不启用动态格式DynamicRankSupportFlag(true)支持动态 rank任意维度数DynamicShapeSupportFlag(true)支持动态 shapeNeedCheckSupportFlag(false)ExtendCfgInfo(opFile.value, invert_apt)指定 kernel 实现文件为 op_kernel/invert_apt.cpp。3.2 输入输出类型一致性校验Tiling 阶段会对输入输出 dtype 做一致性检查见 op_host/arch35/invert_tiling_arch35.cppge::graphStatus InvertTiling::CheckAndGetOutputDtype(ge::DataType outputDtype) { ... OP_CHECK_IF((inputDtype ! outputDtype), ..., return ge::GRAPH_FAILED); ... }若y的 dtype 与x不一致或者落入{DT_INT8, DT_INT16, DT_INT32, DT_INT64, DT_UINT8, DT_UINT16, DT_UINT32, DT_UINT64}之外如浮点、BF16Tiling 会直接报错返回GRAPH_FAILED从而保证算子只按位运算语义合法。四、算子实现架构与调用链路Invert 算子遵循 CANN 算子标准的五段式实现结构目录见 math/invert目录职责关键文件op_graph算子 IRProto定义与图推理invert_proto.h、invert_graph_infer.cppop_host算子定义注册、Infershape、Tilinginvert_def.cpp、invert_infershape.cpp、arch35/invert_tiling_arch35.cppop_kernelNPU 内核实现invert_apt.cpp、arch35/invert_dag.hop_apiL0 API 层aclnn 执行入口invert.cpp、invert.hframework第三方框架TensorFlow算子映射invert_tf_plugin.cpp4.1 算子 IR 定义op_graphinvert_proto.h 使用 GE 的REG_OP宏声明算子接口输入输出类型统一使用TensorType::IntegerDataType()注释中明确其与 TensorFlow 的 Invert 算子兼容REG_OP(Invert) .INPUT(x, TensorType::IntegerDataType()) .OUTPUT(y, TensorType::IntegerDataType()) .OP_END_FACTORY_REG(Invert)同时图推理graph infer逻辑由 invert_graph_infer.cpp 承载用于在构图阶段推导输出 shape。4.2 Infershape复用逐元素通用推理op_host/invert_infershape.cpp 直接复用了逐元素算子的通用 shape 推理IMPL_OP_INFERSHAPE(Invert).InferShape(Ops::Base::InferShape4Elewise);InferShape4Elewise是 CANN 为逐元素算子提供的通用推理工具来源于infershape_elewise_util.hInvert 作为标准 elewise 算子无需单独编写推理逻辑。4.3 Tiling按 dtype 分发到不同 DAG 模板Tiling 是 Ascend 算子在 Host 侧完成的分块调度计算决定 NPU 上如何切分数据、配置核数。Invert 的 Tiling 逻辑arch35/invert_tiling_arch35.cpp要点如下先做输入输出 dtype 一致性检查见 3.2 节根据输出 dtype 分发到对应类型的InvertDagT::OpDag模板调用通用的DoTiling完成 elewise 分块计算SetTilingData()中设置TilingKey 101INVERT_TILING_KEYkernel 侧用TILING_KEY_IS(101UL)匹配workspace 预留 16MBINVERT_WORKSPACE_RESERVE_BYTE 16777216blockNum核数通过context_-SetBlockDim(td_-blockNum)下发。4.4 KernelDAG 描述的 Vector Not 计算Kernel 侧使用DAG 描述 ElementwiseSch 调度器的现代写法。计算核心在 arch35/invert_dag.htemplate typename U struct InvertDag { using OpCopyIn BindVec::CopyInU, Placeholder::In0U; using OpResult BindVec::NotU, OpCopyIn; // 按位取反核心运算 using OpCopyOut BindVec::CopyOutU, Placeholder::Out0U, OpResult; using Outputs ElemsOpCopyOut; using MemCfg MemOptCfgMemLevel::LEVEL_2; // L2 缓存优化 using OpDag DAGSchOutputs, void, MemCfg; };即CopyIn搬运输入→Vec::Not向量按位取反→CopyOut搬出结果并启用 L2 内存级优化。Kernel 入口 op_kernel/invert_apt.cpp 通过if constexpr对 8 种整数类型分别实例化ElementwiseSch使用KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)指定 AIV 向量核执行extern C __global__ __aicore__ void invert(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); REGISTER_TILING_DEFAULT(EleBaseTilingData); ... if (TILING_KEY_IS(101UL)) { if constexpr (std::is_sameDTYPE_X, int8_t::value) { ElementwiseSch0UL, InvertDagint8_t::OpDag sch(tilingData, pipe); sch.Init(x, y); sch.Process(); } // ... 其余 7 种整数类型同理 } }4.5 L0 API 层AICore / AICpu 双路径路由op_api/invert.cpp 是算子对外的高层执行入口实现双路径路由const aclTensor* Invert(const aclTensor* self, aclOpExecutor* executor) { auto out executor-AllocTensor(self-GetViewShape(), self-GetDataType()); if (IsAiCoreSupport(self)) { return InvertAiCore(self, out, executor); // 走 AICore 向量核 } else { return InvertAiCpu(self, out, executor); // 回退 AICputf_kernel } }InvertAiCore通过ADD_TO_LAUNCHER_LIST_AICORE(Invert, OP_INPUT(self), OP_OUTPUT(out))把算子加入 AICore 任务队列InvertAiCpu通过ADD_TO_LAUNCHER_LIST_AICPU(Invert, OP_ATTR_NAMES(), OP_INPUT(self), OP_OUTPUT(out))走 AICputf_kernel执行。RegBase 架构A2/A3 等下 8 种整数类型全部命中 AICore其余架构仅 INT16/UINT16 命中 AICore其余类型回退 AICpu。4.6 TensorFlow 框架适配framework/invert_tf_plugin.cpp 注册了 TF 框架的自定义算子映射REGISTER_CUSTOM_OP(Invert) .FrameworkType(TENSORFLOW) .OriginOpType(Invert) .ParseParamsByOperatorFn(AutoMappingByOpFn) .ImplyType(ImplyType::TVM);这与 invert_proto.h 注释中Compatible with the TensorFlow operator Invert的声明互为印证说明 CANN 的 Invert 与 TensorFlowtf.bitwise.invert语义对齐。五、调用说明与图模式实战README 给出的调用方式如下调用方式调用样例说明图模式调用test_geir_invert.cpp通过算子 IR 构图方式调用 invert 算子5.1 图模式调用完整流程以 examples/test_geir_invert.cpp 为样板图模式调用的完整流程为初始化 GE设置全局选项deviceId、graphRunMode并调用ge::GEInitializestd::mapAscendString, AscendString global_options {{ge.exec.deviceId, 0}, {ge.graphRunMode, 1}}; Status ret ge::GEInitialize(global_options);构造单算子图通过op::Invert(invert1)创建算子节点用ADD_INPUT/ADD_OUTPUT宏绑定输入输出 TensorDescauto invert1 op::Invert(invert1); std::vectorint64_t xShape {4, 2}; ADD_INPUT(1, x, inDtype, xShape); // 生成 placeholder 输入填充值 2 ADD_OUTPUT(1, y, inDtype, xShape); // 声明输出 shape/dtype outputs.push_back(invert1);其中ADD_INPUT宏内部会调用GenOnesData生成 shape 为{4, 2}、元素值为 2 的输入张量示例默认使用DT_INT32值为 2 的 INT32 输入经按位取反后输出 -3。构建并提交图graph.SetInputs(inputs).SetOutputs(outputs)创建Sessionsession-AddGraph(graph_id, graph, graph_options)。Dump 并运行aclgrphDumpGraph(graph, file_path, len)可把构图 dump 到本地默认./dump随后session-RunGraph(graph_id, input, output)执行。回读结果将输入输出数据写为二进制文件命名形如tc_ge_irrun_test_0008_npu_input_0.bin/_output_0.bin并逐元素打印结果验证for (int64_t j 0; j output_shape; j) { LOG_PRINT(result[%ld] is: %d\n, j, result[j]); }收尾ge::GEGetErrorMsgV2()/GEGetWarningMsgV2()获取错误与告警信息最后GEFinalize()释放 GE 会话。5.2 单测覆盖UT 测试 tests/ut/op_api/test_aclnn_invert.cpp 覆盖了多类型用例Invert_int16、Invert_uint16AICore 路径、Invert_int8、Invert_int32等。测试通过aclCreateTensor构造 ND 格式 tensor调用l0op::Invert后断言输出 shape 与输入一致EXPECT_EQ(ToShapeVector(result-GetViewShape()), expectShape)验证了算子输出 shape 继承输入 shape 的逐元素语义。六、总结Invert 是 ops-math 中最典型的逐元素整数算子之一语义单一out ~input支持 8 种整数类型、ND 格式输出与输入 shape、dtype 完全一致实现范式标准Protoinvert_proto.h→ Def/Infershapeinvert_def.cpp、invert_infershape.cpp→ Tilinginvert_tiling_arch35.cpp→ DAG Kernelinvert_dag.h invert_apt.cpp→ L0 APIinvert.cpp架构兼容AICore 向量核Vec::NotElementwiseSch L2 优化为主路径AICpu 兜底RegBase 架构下 8 种整数类型全覆盖调用灵活既可通过图模式GEIR 构图见 test_geir_invert.cpp调用也可通过 L0 API 单算子执行并支持 TensorFlow 框架算子映射。如需在自定义模型中复用 Invert 进行掩码翻转或编码取反可直接参照本文的参数表与图模式样例构造算子图若需在昇腾上做性能优化可进一步阅读 Tiling 与 DAG 源码理解其 L2 缓存与向量化调度策略。赞分享算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载相关推荐CANN ops-math BroadcastTo 算子深度解析NPU 广播实现、图模式调用与源码架构CANN ops math BroadcastTo 算子深度解析NPU 广播实现、图模式调用与源码架构 本文以 CANN ops math 开源算子库中 Br算子库人工智能CANNCANN ops-math 算子深度解析ConcatDV2 多 Tensor 拼接算子的 Tiling 实现与 aclnnCat 调用指南CANN ops math 算子深度解析ConcatDV2 多 Tensor 拼接算子的 Tiling 实现与 aclnnCat 调用指南 ConcatDV2算子库人工智能CANNRightShift 算子深度解析CANN ops-math 按位右移实现与 aclnn 两段式调用指南RightShift 算子深度解析CANN ops math 按位右移实现与 aclnn 两段式调用指南 导读 本文以 CANN ops math 仓库中 R算子库人工智能CANN创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN ops-math Invert 算子深度解析：按位取反的架构实现、Tiling 调度与图模式调用实战
算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载导读Invert 是 CANN ops-math 数学算子库中的按位取反算子对输入张量的每一个元素执行按位取反bitwise NOT运算等价于out_i ~input_i。本文以 math/invert/README.md 为核心骨架结合算子注册、构图、Tiling、Kernel 与单测等源码完整梳理 Invert 算子在 Ascend NPU 上的支持范围、参数语义、底层实现原理与图模式调用方法。读完本文你将掌握Invert 算子的数据类型/产品支持矩阵、逐元素算子的标准实现范式Proto → Def → Infershape → Tiling → Kernel、以及如何用算子 IR 构图方式在 GE 会话中跑通单算子图。一、算子功能与数学定义Invert 是典型的逐元素elementwise / elewise一元算子功能定义非常简洁算子功能对输入张量x的每一个元素执行按位取反bitwise invert得到输出张量y。计算公式$$ out_i\sim input_i $$即每个输出元素是输入元素按位取反后的结果。以 INT8 为例输入0b000000011经 invert 后输出0b11111110-2输入0b11110000则输出0b00001111。由于是按位运算invert(invert(x)) x恒成立因此该算子在符号/掩码翻转、编码反转等场景中常被用作基础原语。需要特别说明的是这里的取反是按位取反bitwise NOT不是算术取负negation-x。算术取负在 CANN 中由独立的 neg 算子承担见 math/neg/README.md两者语义不同~x在补码表示下等价于-x - 1。二、产品支持情况Invert 算子在各产品形态上的支持情况如下表取自 math/invert/README.md产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品√Atlas 训练系列产品√从源码可以交叉印证该支持矩阵。在 op_host/invert_def.cpp 中算子定义通过this-AICore().AddConfig(ascend950, aicoreConfig)显式注册了 ascend950 架构对应 Ascend 950 系列的 AICore 配置在 op_api/invert.cpp 中算子运行时会根据当前 NPU 架构动态选择执行路径// math/invert/op_api/invert.cpp static inline bool IsAiCoreSupport(const aclTensor* self) { auto npuArch op::GetCurrentPlatformInfo().GetCurNpuArch(); if (IsRegBase(npuArch)) { return CheckType(self-GetDataType(), REGBASE_DTYPE_SUPPORT_LIST); } return CheckType(self-GetDataType(), AICORE_DTYPE_SUPPORT_LIST); }即RegBase 类架构如 A2/A3 等通用寄存器基架构支持 8 种整数类型全部走 AICore其余架构仅 INT16、UINT16 两种类型走 AICore其余类型回退到 AICpu 执行。这也解释了为什么 README 中标注 Atlas 200I/500 A2 推理产品不支持——该形态无对应的 AICore/AICpu 内核路径。注意README 中不支持 BFLOAT16的结论在参数表中已明确体现Invert 仅接受 8 种整数类型浮点与 BF16 均不在支持范围内。三、参数说明Invert 算子只有一对输入输出参数取自 math/invert/README.md参数名输入/输出/属性描述数据类型数据格式x输入待进行 invert 计算的入参公式中的 input_iINT8、INT16、INT32、INT64、UINT8、UINT16、UINT32、UINT64NDy输出待进行 invert 计算的出参公式中的 out_iINT8、INT16、INT32、INT64、UINT8、UINT16、UINT32、UINT64ND附加约束Atlas 训练系列产品、Atlas 推理系列产品不支持 BFLOAT16。约束说明无除数据类型约束外Invert 不引入其他 shape、format 层面的限制。3.1 参数语义在算子定义中的落地参数声明与 README 完全对应见 op_host/invert_def.cppthis-Input(x) .ParamType(REQUIRED) .DataType({ge::DT_INT8, ge::DT_INT16, ge::DT_INT32, ge::DT_INT64, ge::DT_UINT8, ge::DT_UINT16, ge::DT_UINT32, ge::DT_UINT64}) .Format({ge::FORMAT_ND, ...}); this-Output(y) .ParamType(REQUIRED) .DataType({...}) // 与 x 完全一致 .Format({ge::FORMAT_ND, ...});该定义同时设置了 AICore 配置项DynamicCompileStaticFlag(true)支持动态编译静态化DynamicFormatFlag(false)不启用动态格式DynamicRankSupportFlag(true)支持动态 rank任意维度数DynamicShapeSupportFlag(true)支持动态 shapeNeedCheckSupportFlag(false)ExtendCfgInfo(opFile.value, invert_apt)指定 kernel 实现文件为 op_kernel/invert_apt.cpp。3.2 输入输出类型一致性校验Tiling 阶段会对输入输出 dtype 做一致性检查见 op_host/arch35/invert_tiling_arch35.cppge::graphStatus InvertTiling::CheckAndGetOutputDtype(ge::DataType outputDtype) { ... OP_CHECK_IF((inputDtype ! outputDtype), ..., return ge::GRAPH_FAILED); ... }若y的 dtype 与x不一致或者落入{DT_INT8, DT_INT16, DT_INT32, DT_INT64, DT_UINT8, DT_UINT16, DT_UINT32, DT_UINT64}之外如浮点、BF16Tiling 会直接报错返回GRAPH_FAILED从而保证算子只按位运算语义合法。四、算子实现架构与调用链路Invert 算子遵循 CANN 算子标准的五段式实现结构目录见 math/invert目录职责关键文件op_graph算子 IRProto定义与图推理invert_proto.h、invert_graph_infer.cppop_host算子定义注册、Infershape、Tilinginvert_def.cpp、invert_infershape.cpp、arch35/invert_tiling_arch35.cppop_kernelNPU 内核实现invert_apt.cpp、arch35/invert_dag.hop_apiL0 API 层aclnn 执行入口invert.cpp、invert.hframework第三方框架TensorFlow算子映射invert_tf_plugin.cpp4.1 算子 IR 定义op_graphinvert_proto.h 使用 GE 的REG_OP宏声明算子接口输入输出类型统一使用TensorType::IntegerDataType()注释中明确其与 TensorFlow 的 Invert 算子兼容REG_OP(Invert) .INPUT(x, TensorType::IntegerDataType()) .OUTPUT(y, TensorType::IntegerDataType()) .OP_END_FACTORY_REG(Invert)同时图推理graph infer逻辑由 invert_graph_infer.cpp 承载用于在构图阶段推导输出 shape。4.2 Infershape复用逐元素通用推理op_host/invert_infershape.cpp 直接复用了逐元素算子的通用 shape 推理IMPL_OP_INFERSHAPE(Invert).InferShape(Ops::Base::InferShape4Elewise);InferShape4Elewise是 CANN 为逐元素算子提供的通用推理工具来源于infershape_elewise_util.hInvert 作为标准 elewise 算子无需单独编写推理逻辑。4.3 Tiling按 dtype 分发到不同 DAG 模板Tiling 是 Ascend 算子在 Host 侧完成的分块调度计算决定 NPU 上如何切分数据、配置核数。Invert 的 Tiling 逻辑arch35/invert_tiling_arch35.cpp要点如下先做输入输出 dtype 一致性检查见 3.2 节根据输出 dtype 分发到对应类型的InvertDagT::OpDag模板调用通用的DoTiling完成 elewise 分块计算SetTilingData()中设置TilingKey 101INVERT_TILING_KEYkernel 侧用TILING_KEY_IS(101UL)匹配workspace 预留 16MBINVERT_WORKSPACE_RESERVE_BYTE 16777216blockNum核数通过context_-SetBlockDim(td_-blockNum)下发。4.4 KernelDAG 描述的 Vector Not 计算Kernel 侧使用DAG 描述 ElementwiseSch 调度器的现代写法。计算核心在 arch35/invert_dag.htemplate typename U struct InvertDag { using OpCopyIn BindVec::CopyInU, Placeholder::In0U; using OpResult BindVec::NotU, OpCopyIn; // 按位取反核心运算 using OpCopyOut BindVec::CopyOutU, Placeholder::Out0U, OpResult; using Outputs ElemsOpCopyOut; using MemCfg MemOptCfgMemLevel::LEVEL_2; // L2 缓存优化 using OpDag DAGSchOutputs, void, MemCfg; };即CopyIn搬运输入→Vec::Not向量按位取反→CopyOut搬出结果并启用 L2 内存级优化。Kernel 入口 op_kernel/invert_apt.cpp 通过if constexpr对 8 种整数类型分别实例化ElementwiseSch使用KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)指定 AIV 向量核执行extern C __global__ __aicore__ void invert(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); REGISTER_TILING_DEFAULT(EleBaseTilingData); ... if (TILING_KEY_IS(101UL)) { if constexpr (std::is_sameDTYPE_X, int8_t::value) { ElementwiseSch0UL, InvertDagint8_t::OpDag sch(tilingData, pipe); sch.Init(x, y); sch.Process(); } // ... 其余 7 种整数类型同理 } }4.5 L0 API 层AICore / AICpu 双路径路由op_api/invert.cpp 是算子对外的高层执行入口实现双路径路由const aclTensor* Invert(const aclTensor* self, aclOpExecutor* executor) { auto out executor-AllocTensor(self-GetViewShape(), self-GetDataType()); if (IsAiCoreSupport(self)) { return InvertAiCore(self, out, executor); // 走 AICore 向量核 } else { return InvertAiCpu(self, out, executor); // 回退 AICputf_kernel } }InvertAiCore通过ADD_TO_LAUNCHER_LIST_AICORE(Invert, OP_INPUT(self), OP_OUTPUT(out))把算子加入 AICore 任务队列InvertAiCpu通过ADD_TO_LAUNCHER_LIST_AICPU(Invert, OP_ATTR_NAMES(), OP_INPUT(self), OP_OUTPUT(out))走 AICputf_kernel执行。RegBase 架构A2/A3 等下 8 种整数类型全部命中 AICore其余架构仅 INT16/UINT16 命中 AICore其余类型回退 AICpu。4.6 TensorFlow 框架适配framework/invert_tf_plugin.cpp 注册了 TF 框架的自定义算子映射REGISTER_CUSTOM_OP(Invert) .FrameworkType(TENSORFLOW) .OriginOpType(Invert) .ParseParamsByOperatorFn(AutoMappingByOpFn) .ImplyType(ImplyType::TVM);这与 invert_proto.h 注释中Compatible with the TensorFlow operator Invert的声明互为印证说明 CANN 的 Invert 与 TensorFlowtf.bitwise.invert语义对齐。五、调用说明与图模式实战README 给出的调用方式如下调用方式调用样例说明图模式调用test_geir_invert.cpp通过算子 IR 构图方式调用 invert 算子5.1 图模式调用完整流程以 examples/test_geir_invert.cpp 为样板图模式调用的完整流程为初始化 GE设置全局选项deviceId、graphRunMode并调用ge::GEInitializestd::mapAscendString, AscendString global_options {{ge.exec.deviceId, 0}, {ge.graphRunMode, 1}}; Status ret ge::GEInitialize(global_options);构造单算子图通过op::Invert(invert1)创建算子节点用ADD_INPUT/ADD_OUTPUT宏绑定输入输出 TensorDescauto invert1 op::Invert(invert1); std::vectorint64_t xShape {4, 2}; ADD_INPUT(1, x, inDtype, xShape); // 生成 placeholder 输入填充值 2 ADD_OUTPUT(1, y, inDtype, xShape); // 声明输出 shape/dtype outputs.push_back(invert1);其中ADD_INPUT宏内部会调用GenOnesData生成 shape 为{4, 2}、元素值为 2 的输入张量示例默认使用DT_INT32值为 2 的 INT32 输入经按位取反后输出 -3。构建并提交图graph.SetInputs(inputs).SetOutputs(outputs)创建Sessionsession-AddGraph(graph_id, graph, graph_options)。Dump 并运行aclgrphDumpGraph(graph, file_path, len)可把构图 dump 到本地默认./dump随后session-RunGraph(graph_id, input, output)执行。回读结果将输入输出数据写为二进制文件命名形如tc_ge_irrun_test_0008_npu_input_0.bin/_output_0.bin并逐元素打印结果验证for (int64_t j 0; j output_shape; j) { LOG_PRINT(result[%ld] is: %d\n, j, result[j]); }收尾ge::GEGetErrorMsgV2()/GEGetWarningMsgV2()获取错误与告警信息最后GEFinalize()释放 GE 会话。5.2 单测覆盖UT 测试 tests/ut/op_api/test_aclnn_invert.cpp 覆盖了多类型用例Invert_int16、Invert_uint16AICore 路径、Invert_int8、Invert_int32等。测试通过aclCreateTensor构造 ND 格式 tensor调用l0op::Invert后断言输出 shape 与输入一致EXPECT_EQ(ToShapeVector(result-GetViewShape()), expectShape)验证了算子输出 shape 继承输入 shape 的逐元素语义。六、总结Invert 是 ops-math 中最典型的逐元素整数算子之一语义单一out ~input支持 8 种整数类型、ND 格式输出与输入 shape、dtype 完全一致实现范式标准Protoinvert_proto.h→ Def/Infershapeinvert_def.cpp、invert_infershape.cpp→ Tilinginvert_tiling_arch35.cpp→ DAG Kernelinvert_dag.h invert_apt.cpp→ L0 APIinvert.cpp架构兼容AICore 向量核Vec::NotElementwiseSch L2 优化为主路径AICpu 兜底RegBase 架构下 8 种整数类型全覆盖调用灵活既可通过图模式GEIR 构图见 test_geir_invert.cpp调用也可通过 L0 API 单算子执行并支持 TensorFlow 框架算子映射。如需在自定义模型中复用 Invert 进行掩码翻转或编码取反可直接参照本文的参数表与图模式样例构造算子图若需在昇腾上做性能优化可进一步阅读 Tiling 与 DAG 源码理解其 L2 缓存与向量化调度策略。赞分享算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载相关推荐CANN ops-math BroadcastTo 算子深度解析NPU 广播实现、图模式调用与源码架构CANN ops math BroadcastTo 算子深度解析NPU 广播实现、图模式调用与源码架构 本文以 CANN ops math 开源算子库中 Br算子库人工智能CANNCANN ops-math 算子深度解析ConcatDV2 多 Tensor 拼接算子的 Tiling 实现与 aclnnCat 调用指南CANN ops math 算子深度解析ConcatDV2 多 Tensor 拼接算子的 Tiling 实现与 aclnnCat 调用指南 ConcatDV2算子库人工智能CANNRightShift 算子深度解析CANN ops-math 按位右移实现与 aclnn 两段式调用指南RightShift 算子深度解析CANN ops math 按位右移实现与 aclnn 两段式调用指南 导读 本文以 CANN ops math 仓库中 R算子库人工智能CANN创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考