人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载导读本文基于 CANN 神经网络算子库 ops-nn 中 foreach_binary_op 算子文档系统讲解 ForeachBinaryOp 这一图融合fused图内部算子的功能定义、算子原型、Tiling 调度原理、SIMT Kernel 实现与 GE IR 构图调用方式。该算子面向 Ascend 950arch35产品将 add/sub/mul/div 四类逐 Tensor 列表二元运算统一为一个算子供图模式GE IR / 图融合 Pass内部使用。读完本文你将掌握该算子的输入输出约束、op_code与 TilingKey 的映射规则、多 Tensor 跨核切分机制以及如何通过 GE IR 接口在图中正确构图调用它。前置说明ForeachBinaryOp不对外提供 aclnn 单算子接口只能以图内部算子的形式出现在 GE IR 图中由上层框架的图融合 Pass 生成最终在 Ascend 950 上以 SIMT Kernel 方式执行。一、产品支持情况该算子为 arch35Ascend 950专用算子产品支持矩阵如下产品是否支持Ascend 950PR / Ascend 950DTarch35/ascend950√其它产品×这一结论在算子注册代码中有直接印证foreach_binary_op_def.cpp 中通过OpAICoreConfig仅调用了AddConfig(ascend950, aicoreConfig950)即该算子只在 ascend950 平台上注册了 AICore 配置Tiling 与 Kernel 实现也都位于arch35目录下。二、功能说明与计算公式2.1 功能定位ForeachBinaryOp 对两个 Tensor 列表x1、x2逐 Tensor、逐元素做二元运算运算类型由编译期属性op_code选择0add、1sub、2mul、3div。将多种二元 foreach 运算统一为一个算子便于图融合 Pass 将逐个列表元素分别执行二元算子的图模式折叠成单个算子从而降低构图开销、提升调度效率。计算公式为$$ x1 [{x1_0}, {x1_1}, ... {x1_{n-1}}],\ x2 [{x2_0}, {x2_1}, ... {x2_{n-1}}],\ y [{y_0}, {y_1}, ... {y_{n-1}}] $$$$ y_i x1_i \odot x2_i \quad (i0,1,...,n-1) $$其中 $\odot$ 由op_code决定op_code运算公式0add$y_i x1_i x2_i$1sub$y_i x1_i - x2_i$2mul$y_i x1_i \times x2_i$3div$y_i x1_i / x2_i$2.2 除零语义的特殊处理文档特别指出整型INT32除法对除数为 0 的元素结果置 0以规避设备上整型除零的未定义行为可能触发 trap而浮点FP16/FP32/BF16除以 0 时遵循 IEEE 语义产生 inf/nan不做额外处理。这一语义在 Kernel 源码 foreach_binary_op_simt.h 的BinaryApply模板函数中有精确实现if constexpr (OP OP_CODE_ADD) { return a b; } else if constexpr (OP OP_CODE_SUB) { return a - b; } else if constexpr (OP OP_CODE_MUL) { return a * b; } else { // Integer divide-by-zero is undefined on device (may trap); guard it. Float div by 0 // yields IEEE inf/nan which is well-defined and left as-is. if constexpr (std::is_integral_vT) { return (b static_castT(0)) ? static_castT(0) : (a / b); } else { return a / b; } }可以看到BinaryApply是一个编译期模板分派函数OP作为模板参数在编译期决定分支if constexpr避免运行时跳转除零保护仅对整型生效浮点路径直接返回a / b。三、算子定义REG_OP 原型算子原型注册于 foreach_binary_op_proto.h使用REG_OP(ForeachBinaryOp)宏声明两个动态输入、一个动态输出与一个必选属性REG_OP(ForeachBinaryOp) .DYNAMIC_INPUT(x1, TensorType({DT_FLOAT, DT_FLOAT16, DT_INT32, DT_BF16})) .DYNAMIC_INPUT(x2, TensorType({DT_FLOAT, DT_FLOAT16, DT_INT32, DT_BF16})) .DYNAMIC_OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16, DT_INT32, DT_BF16})) .REQUIRED_ATTR(op_code, Int) .OP_END_FACTORY_REG(ForeachBinaryOp)3.1 输入参数名类型描述数据类型数据格式x1DYNAMIC_INPUTTensorList第一个输入张量列表对应公式中的x1。列表内所有 Tensor 的数据类型一致。FLOAT16、FLOAT、INT32、BFLOAT16NDx2DYNAMIC_INPUTTensorList第二个输入张量列表对应公式中的x2。数据类型、shape 与x1一致。FLOAT16、FLOAT、INT32、BFLOAT16ND3.2 属性属性名类型必选描述op_codeInt是REQUIRED_ATTR选择二元运算0add、1sub、2mul、3div。取值范围 [0, 4)。3.3 输出参数名类型描述数据类型数据格式yDYNAMIC_OUTPUTTensorList输出张量列表对应公式中的yy_i x1_i op x2_i。数据类型、shape 与x1一致。FLOAT16、FLOAT、INT32、BFLOAT16ND3.4 host 侧算子定义补充除 GE 图 IR 原型外host 侧还通过OpDef机制补充了算子在编译期的完整描述见 foreach_binary_op_def.cpp。其中值得注意的配置项Input(x1)/Input(x2)/Output(y)均为ParamType(DYNAMIC)数据类型列表与 proto 一致格式固定为FORMAT_NDAutoContiguous()声明输入为连续内存布局供框架优化访存DynamicCompileStaticFlag(true)、DynamicShapeSupportFlag(true)、DynamicRankSupportFlag(true)表明该算子支持动态 shape 与动态 rankPrecisionReduceFlag(true)允许精度相关优化ExtendCfgInfo(opFile.value, foreach_binary_op)将算子实现关联到名为foreach_binary_op的 kernel 二进制文件。四、约束说明仅支持 Ascend 950PR/Ascend 950DTarch35/ascend950SIMT Kernel 实现。x1、x2、y三个列表长度Tensor 个数一致且一一对应的 Tensor shape 一致。列表 Tensor 个数上限为 256MAX_TENSOR_NUM_FOREACH_BINARY_OP。x1、x2、y中每个 Tensor 的数据类型一致且属于 FLOAT16/FLOAT/INT32/BFLOAT16。支持空 Tensor总元素数为 0 时不启核。op_code必须落在 [0, 4) 区间否则 Tiling 报错。TilingKey 规则tilingKey op_code * 4 dtypeIdx其中 dtypeIdxFP160、FP321、INT322、BF163共 16 种调度模式 [0, 15]。其中列表 Tensor 个数上限 256在 host 与 kernel 两侧都有常量定义host 侧见 foreach_binary_op_tiling_arch35.hMAX_TENSOR_NUM_FOREACH_BINARY_OP 256kernel 侧见 foreach_binary_op_tiling_data.hMAX_TENSOR_CONT 256同时MAX_CORE_CONT 80与 ascend950 的 AIV 核数对应。五、TilingKey 调度机制ForeachBinaryOp 的调度核心是一码定模式的 TilingKey 设计tilingKey op_code * 4 dtypeIdx其中 dtypeIdx 由输入数据类型映射而来FP160、FP321、INT322、BF163共覆盖 4 种运算 × 4 种 dtype 16 种调度模式。host 侧在 foreach_binary_op_tiling_arch35.cpp 的GetDtypeIdx中完成 dtype→索引映射并在第 194 行计算tilingKey opCode * DTYPE_NUM GetDtypeIdx(dataType)kernel 侧在 foreach_binary_op_tiling_key.h 通过ASCENDC_TPL_ARGS_DECL/ASCENDC_TPL_SEL将 0~15 共 16 个整型枚举为编译期模板参数schMode即每个 TilingKey 对应一份模板特化实例化算子入口 foreach_binary_op.cpp 在编译期从schMode反解出opCode schMode / 4与dtypeIdx schMode % 4再通过if constexpr选择对应的数据类型处理路径constexpr uint32_t opCode schMode / 4; constexpr uint32_t dtypeIdx schMode % 4; if constexpr (dtypeIdx 0) { ForeachBinaryOpProcessFp16opCode(x1, x2, y, tilingData); } else if constexpr (dtypeIdx 1) { ForeachBinaryOpProcessFp32opCode(x1, x2, y, tilingData); } else if constexpr (dtypeIdx 2) { ForeachBinaryOpProcessInt32opCode(x1, x2, y, tilingData); } else { ForeachBinaryOpProcessBf16opCode(x1, x2, y, tilingData); }由于 op_code 与 dtype 全部在编译期确定Kernel 内不会出现任何与运算类型或数据类型相关的运行时分支这是该算子实现高性能的关键设计之一。六、Tiling 原理多 Tensor 跨核切分Tiling 函数实现在 foreach_binary_op_tiling_arch35.cpp整体流程如下获取平台信息GetCoreAndUbSize优先从编译期ForeachBinaryOpCompileInfocoreNum/ubSize读取否则回退到PlatformAscendC获取 AIV 核数与 UB 大小UB 需先扣除DCACHE_SIZE 128 * 1024字节的 data cache 预留后才是 Kernel 可用的本地内存大小。校验 op_codeGetOpCode从属性中读取op_code若不在 [0, 4) 区间内直接返回GRAPH_FAILED与文档tiling 报错的约束一致。统计 Tensor 个数GetTensorNum通过GetInputInstanceInfo(0)-GetInstanceNum()获取x1列表长度并校验不超过MAX_TENSOR_NUM_FOREACH_BINARY_OP256。填充逐 Tensor 元素数FillTensorDataCount遍历每个输入 Tensor 的StorageShape将各自元素数写入tensorDataCountList[i]并累加totalElements首个 Tensor 的 dtype 被记录为该次调度的数据类型。空 Tensor 短路当totalElements 0时needCoreNum置 0Kernel 入口检测到needCoreNum 0直接 return即不启核但SetBlockDim(1)仍保证图调度合法TilingKey 照常下发。计算核数与每核元素数needCoreNum min(coreNum, ceil(totalElements / SINGLE_CORE_MIN_ELEMENTS))即每核至少承担 1024 个元素避免小任务过度并行perCoreElements align_up(ceil(totalElements / needCoreNum), 32)按 32 对齐。跨 Tensor 分配AssignCoresToTensors该函数定义在头文件 foreach_binary_op_tiling_arch35.h 中内联在头文件内以便 UT 白盒测试。其核心思想是以全局元素序号为单位切分数据再回扫定位每个核覆盖的 Tensor 区间。对每个核记录tensorStartList[core]/tensorStartOffsetList[core]起始 Tensor 下标及该 Tensor 内的起始偏移tensorEndList[core]/tensorEndOffsetList[core]结束 Tensor 下标及该 Tensor 内的结束偏移含切分按 32 对齐向上取整后前序核可能覆盖全部元素剩余核会因coreStart totalElements提前 break最终返回实际启用的核数usedCoreNum确保SetBlockDim不会启动没有分配区间的空核。收尾FinalizeTilingSetBlockDim(needCoreNum)、SetTilingKey(tilingKey)、SetLocalMemorySize(ubSize)workspace 申请清零。对应的 Tiling 数据宿主结构见 foreach_binary_op_tiling_data.hKernel 侧与 host 侧字段一一对应needCoreNum、tensorCount、totalElements、tensorDataCountList[256]、tensorStartList[80]等其中 80 对应 ascend950 的 AIV 核数上限。七、SIMT Kernel 实现细节Kernel 实现在 foreach_binary_op_simt.h采用 SIMTSIMT VF编程模型针对四种数据类型提供了三条计算路径数据类型实现方式说明FP32直接计算OpForeachBinaryDirectSimtfloat精度直接满足要求INT32直接计算OpForeachBinaryDirectSimtint32_t直接整型运算FP16提升为 FP32 中间计算OpForeachBinaryFp16CastSimt__half2float转 FP32 计算后再__float2half_rn回写保证精度BF16提升为 FP32 中间计算OpForeachBinaryBf16CastSimt__bfloat162float转 FP32 计算后再__float2bfloat16_rn回写7.1 线程组织与索引位宽分派线程数由索引类型位宽决定THREAD_NUM_VF (sizeof(IDX_T) 4) ? 1024 : 512即 32 位索引用 1024 线程、64 位索引用 512 线程每个核的局部元素数localCount若不超过INT32_MAX则使用int32_t索引调用 VF 核否则使用int64_t索引从而在核内分段元素数极大21 亿时仍能正确索引。7.2 多 Tensor 遍历与 ListTensorDesc每个处理函数如ForeachBinaryOpProcessFp32都通过ListTensorDesc读取列表型输入ListTensorDesc x1List(reinterpret_cast__gm__ void*(x1)); ListTensorDesc x2List(reinterpret_cast__gm__ void*(x2)); ListTensorDesc yList(reinterpret_cast__gm__ void*(y)); for (int32_t t startT; t endT; t) { __gm__ float* x1_t x1List.GetDataPtrfloat(t); ... int64_t localStart (t startT) ? tilingData-tensorStartOffsetList[coreId] : 0; int64_t localEnd (t endT) ? tilingData-tensorEndOffsetList[coreId] 1 : totalCount; int64_t localCount localEnd - localStart; if (localCount 0) { ... asc_vf_call...(...) } }即每个核只遍历 Tiling 分配给自己的[startT, endT]Tensor 区间区间首尾 Tensor 用偏移裁剪中间 Tensor 全量处理。随后调用asc_vf_call启动 SIMT VF 核内层for循环以threadIdx.x起步、blockDim.x步长遍历实现元素级并行。八、GE IR 构图调用示例ForeachBinaryOp 是图内部融合算子通过 GE IR 图模式构图调用。完整可运行示例见 test_geir_foreach_binary_op.cpp核心构图片段如下auto op1 op::ForeachBinaryOp(foreachBinaryOp1); const int N 2; // 列表内 Tensor 个数 op1.set_attr_op_code(0); // 0add, 1sub, 2mul, 3div op1.create_dynamic_input_x1(N); for (int i 0; i N; i) { op1.set_dynamic_input_x1(i, dataX1[i]); } op1.create_dynamic_input_x2(N); for (int i 0; i N; i) { op1.set_dynamic_input_x2(i, dataX2[i]); } op1.create_dynamic_output_y(N);示例中完整展示了 GE 图模式的调用流程test_geir_foreach_binary_op.cpp通过op::ForeachBinaryOp(foreachBinaryOp1)创建算子实例命名必须与原型注册名一致set_attr_op_code(0)设置二元运算类型本例为 add编译期即被 Tiling/Kernel 读取create_dynamic_input_x1(N)/create_dynamic_input_x2(N)创建长度 N 的动态输入列表随后逐个set_dynamic_input_x1(i, d)/set_dynamic_input_x2(i, d)绑定数据算子op::Datacreate_dynamic_output_y(N)创建输出列表并逐个update_dynamic_output_desc_y(i, yd)更新输出描述将算子的输入/输出算子注册到Graph通过Session::AddGraph与Session::RunGraph完成构图与执行。示例中每个 Tensor 的 shape 为{256}、dtype 为 FP32MkF32填充 2.0f对应 TilingKey 0 * 4 1 1add FP32。九、测试与验证该算子配套了 host 侧单元测试覆盖文档列出的全部关键约束test_foreach_binary_op_tiling.cpp 的all_op_dtype_combos用例遍历 4 种 op_code × 4 种 dtype 共 16 种组合逐一断言 TilingKey opCode * 4 DtypeIdx(dtype)验证 TilingKey 规则multi_tensor_assign_cores用例白盒测试AssignCoresToTensors16 个 Tensor × 64 元素2 核切分后每个核跨 8 个 Tensor精确断言每个核的起始/结束 Tensor 下标与偏移同时验证核数多于数据需求时尾核 break 并被裁剪的行为empty_no_core_split用例验证空输入shape{0}时 Tiling 仍返回 SUCCESSneedCoreNum 0不启核large_unaligned_count用例用 150004 个元素非 32 的倍数驱动多核切分与 32 对齐路径invalid_op_code用例验证op_code 4越界时 Tiling 返回GRAPH_FAILED与文档op_code 必须落在 [0, 4) 区间否则 tiling 报错完全对应。此外 tests/ut/op_host/arch35/test_foreach_binary_op_infershape.cpp 覆盖动态列表的 Infershape 推导op_graph/fusion_pass目录存放图融合 Pass 侧的配套逻辑op_host/config/ascend950/foreach_binary_op_binary.json为 ascend950 平台的算子二进制配置共同构成proto 定义 → host Tiling → kernel 执行 → 图融合接入的完整闭环。十、总结ForeachBinaryOp 是 CANN ops-nn 中面向 Ascend 950 的一个典型图融合内部算子其设计要点可归纳为一算子多语义用op_code在编译期统一 add/sub/mul/div 四种运算减少图融合后算子种类TilingKey 合一编码op_code * 4 dtypeIdx将运算类型与数据类型编码为 16 种编译期调度模式消除运行时分支跨 Tensor 全局切分Tiling 以全局元素序为单位分配核区间核可横跨多个 Tensor并通过tensorStart/EndList与 offset 精确描述边界精度与健壮性兼顾FP16/BF16 提升到 FP32 中间计算INT32 除零置 0 规避未定义行为空列表不启核严格的使用边界仅 GE IR 图模式可用无 aclnn 接口仅支持 arch35/ascend950列表长度上限 256dtype 限定为 FP16/FP32/INT32/BF16 且格式为 ND。开发者若需在自有框架的图融合 Pass 中生成该算子可参照 test_geir_foreach_binary_op.cpp 的构图方式并严格遵守第四节列出的约束条件。赞分享人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载相关推荐CANN ops-math MemSetV2 算子深度解析Tensor 内存批量初始化与 GE IR 图模式调用实战CANN ops math MemSetV2 算子深度解析Tensor 内存批量初始化与 GE IR 图模式调用实战 MemSetV2 是 CANN ops算子库人工智能CANNCANN ops-nn FusedBiasLeakyReluGrad 算子深度解析从 GE IR 构图到 NPU 反向核函数实现CANN ops nn FusedBiasLeakyReluGrad 算子深度解析从 GE IR 构图到 NPU 反向核函数实现 FusedBiasLeaky人工智能算子库深度学习CANNAscendCANN ops-nn 张量列表混合运算算子 aclnnForeachAddcmulScalarV2 深度解析与实战指南CANN ops nn 张量列表混合运算算子 aclnnForeachAddcmulScalarV2 深度解析与实战指南 aclnnForeachAddcmul人工智能算子库深度学习CANNAscend上一篇Warpgate API集成指南自动化用户与目标管理下一篇spotify-player的编译时配置条件编译与特性标志创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考