企业官网建设流程全解析

• 人工智能
• 算子库
• 深度学习
• CANN
• Ascend

导读

本文基于 CANN 神经网络算子库 ops-nn 中 foreach_binary_op 算子文档，系统讲解 ForeachBinaryOp 这一图融合（fused）图内部算子的功能定义、算子原型、Tiling 调度原理、SIMT Kernel 实现与 GE IR 构图调用方式。该算子面向 Ascend 950（arch35）产品，将 add/sub/mul/div 四类"逐 Tensor 列表二元运算"统一为一个算子，供图模式（GE IR / 图融合 Pass）内部使用。读完本文，你将掌握该算子的输入输出约束、op_code与 TilingKey 的映射规则、多 Tensor 跨核切分机制，以及如何通过 GE IR 接口在图中正确构图调用它。

前置说明：ForeachBinaryOp不对外提供 aclnn 单算子接口，只能以图内部算子的形式出现在 GE IR 图中，由上层框架的图融合 Pass 生成，最终在 Ascend 950 上以 SIMT Kernel 方式执行。

一、产品支持情况

该算子为 arch35（Ascend 950）专用算子，产品支持矩阵如下：

这一结论在算子注册代码中有直接印证：foreach_binary_op_def.cpp 中通过OpAICoreConfig仅调用了AddConfig("ascend950", aicoreConfig950)，即该算子只在 ascend950 平台上注册了 AICore 配置；Tiling 与 Kernel 实现也都位于arch35目录下。

二、功能说明与计算公式

2.1 功能定位

ForeachBinaryOp 对两个 Tensor 列表x1、x2逐 Tensor、逐元素做二元运算，运算类型由编译期属性op_code选择（0=add、1=sub、2=mul、3=div）。将多种二元 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 (i=0,1,...,n-1) $$

其中 $\odot$ 由op_code决定：

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_v<T>) { return (b == static_cast<T>(0)) ? static_cast<T>(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 输入

3.2 属性

3.3 输出

3.4 host 侧算子定义补充

除 GE 图 IR 原型外，host 侧还通过OpDef机制补充了算子在编译期的完整描述，见 foreach_binary_op_def.cpp。其中值得注意的配置项：

• Input("x1")/Input("x2")/Output("y")均为ParamType(DYNAMIC)，数据类型列表与 proto 一致，格式固定为FORMAT_ND；
• AutoContiguous()声明输入为连续内存布局，供框架优化访存；
• DynamicCompileStaticFlag(true)、DynamicShapeSupportFlag(true)、DynamicRankSupportFlag(true)表明该算子支持动态 shape 与动态 rank；
• PrecisionReduceFlag(true)允许精度相关优化；
• ExtendCfgInfo("opFile.value", "foreach_binary_op")将算子实现关联到名为foreach_binary_op的 kernel 二进制文件。

四、约束说明

• 仅支持 Ascend 950PR/Ascend 950DT（arch35/ascend950），SIMT Kernel 实现。
• x1、x2、y三个列表长度（Tensor 个数）一致，且一一对应的 Tensor shape 一致。
• 列表 Tensor 个数上限为 256（MAX_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，其中 dtypeIdx：FP16=0、FP32=1、INT32=2、BF16=3，共 16 种调度模式 [0, 15]。

其中"列表 Tensor 个数上限 256"在 host 与 kernel 两侧都有常量定义：host 侧见 foreach_binary_op_tiling_arch35.h（MAX_TENSOR_NUM_FOREACH_BINARY_OP = 256），kernel 侧见 foreach_binary_op_tiling_data.h（MAX_TENSOR_CONT = 256，同时MAX_CORE_CONT = 80与 ascend950 的 AIV 核数对应）。

五、TilingKey 调度机制

ForeachBinaryOp 的调度核心是"一码定模式"的 TilingKey 设计：tilingKey = op_code * 4 + dtypeIdx，其中 dtypeIdx 由输入数据类型映射而来（FP16=0、FP32=1、INT32=2、BF16=3），共覆盖 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) { ForeachBinaryOpProcessFp16<opCode>(x1, x2, y, &tilingData); } else if constexpr (dtypeIdx == 1) { ForeachBinaryOpProcessFp32<opCode>(x1, x2, y, &tilingData); } else if constexpr (dtypeIdx == 2) { ForeachBinaryOpProcessInt32<opCode>(x1, x2, y, &tilingData); } else { ForeachBinaryOpProcessBf16<opCode>(x1, x2, y, &tilingData); }

由于 op_code 与 dtype 全部在编译期确定，Kernel 内不会出现任何与运算类型或数据类型相关的运行时分支，这是该算子实现高性能的关键设计之一。

六、Tiling 原理：多 Tensor 跨核切分

Tiling 函数实现在 foreach_binary_op_tiling_arch35.cpp，整体流程如下：

1. 获取平台信息（GetCoreAndUbSize）：优先从编译期ForeachBinaryOpCompileInfo（coreNum/ubSize）读取，否则回退到PlatformAscendC获取 AIV 核数与 UB 大小；UB 需先扣除DCACHE_SIZE = 128 * 1024字节的 data cache 预留后，才是 Kernel 可用的本地内存大小。

2. 校验 op_code（GetOpCode）：从属性中读取op_code，若不在 [0, 4) 区间内直接返回GRAPH_FAILED，与文档"tiling 报错"的约束一致。

3. 统计 Tensor 个数（GetTensorNum）：通过GetInputInstanceInfo(0)->GetInstanceNum()获取x1列表长度，并校验不超过MAX_TENSOR_NUM_FOREACH_BINARY_OP（256）。

4. 填充逐 Tensor 元素数（FillTensorDataCount）：遍历每个输入 Tensor 的StorageShape，将各自元素数写入tensorDataCountList[i]并累加totalElements；首个 Tensor 的 dtype 被记录为该次调度的数据类型。

5. 空 Tensor 短路：当totalElements == 0时，needCoreNum置 0（Kernel 入口检测到needCoreNum == 0直接 return，即"不启核"），但SetBlockDim(1)仍保证图调度合法，TilingKey 照常下发。

6. 计算核数与每核元素数：

   • needCoreNum = min(coreNum, ceil(totalElements / SINGLE_CORE_MIN_ELEMENTS))，即每核至少承担 1024 个元素，避免小任务过度并行；
   • perCoreElements = align_up(ceil(totalElements / needCoreNum), 32)，按 32 对齐。

7. 跨 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不会启动没有分配区间的空核。

8. 收尾（FinalizeTiling）：SetBlockDim(needCoreNum)、SetTilingKey(tilingKey)、SetLocalMemorySize(ubSize)，workspace 申请清零。

对应的 Tiling 数据宿主结构见 foreach_binary_op_tiling_data.h，Kernel 侧与 host 侧字段一一对应（needCoreNum、tensorCount、totalElements、tensorDataCountList[256]、tensorStartList[80]等，其中 80 对应 ascend950 的 AIV 核数上限）。

七、SIMT Kernel 实现细节

Kernel 实现在 foreach_binary_op_simt.h，采用 SIMT（SIMT VF）编程模型，针对四种数据类型提供了三条计算路径：

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.GetDataPtr<float>(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); // 0=add, 1=sub, 2=mul, 3=div 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）：

1. 通过op::ForeachBinaryOp("foreachBinaryOp1")创建算子实例，命名必须与原型注册名一致；
2. set_attr_op_code(0)设置二元运算类型（本例为 add），编译期即被 Tiling/Kernel 读取；
3. create_dynamic_input_x1(N)/create_dynamic_input_x2(N)创建长度 N 的动态输入列表，随后逐个set_dynamic_input_x1(i, d)/set_dynamic_input_x2(i, d)绑定数据算子op::Data；
4. create_dynamic_output_y(N)创建输出列表，并逐个update_dynamic_output_desc_y(i, yd)更新输出描述；
5. 将算子的输入/输出算子注册到Graph，通过Session::AddGraph与Session::RunGraph完成构图与执行。

示例中每个 Tensor 的 shape 为{256}、dtype 为 FP32（MkF32填充 2.0f），对应 TilingKey =0 * 4 + 1 = 1（add + 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用例白盒测试AssignCoresToTensors：16 个 Tensor × 64 元素，2 核切分后每个核跨 8 个 Tensor，精确断言每个核的起始/结束 Tensor 下标与偏移；同时验证"核数多于数据需求时，尾核 break 并被裁剪"的行为；
• empty_no_core_split用例验证空输入（shape{0}）时 Tiling 仍返回 SUCCESS，needCoreNum = 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 的一个典型图融合内部算子，其设计要点可归纳为：

1. 一算子多语义：用op_code在编译期统一 add/sub/mul/div 四种运算，减少图融合后算子种类；
2. TilingKey 合一编码：op_code * 4 + dtypeIdx将运算类型与数据类型编码为 16 种编译期调度模式，消除运行时分支；
3. 跨 Tensor 全局切分：Tiling 以全局元素序为单位分配核区间，核可横跨多个 Tensor，并通过tensorStart/EndList与 offset 精确描述边界；
4. 精度与健壮性兼顾：FP16/BF16 提升到 FP32 中间计算，INT32 除零置 0 规避未定义行为，空列表不启核；
5. 严格的使用边界：仅 GE IR 图模式可用（无 aclnn 接口），仅支持 arch35/ascend950，列表长度上限 256，dtype 限定为 FP16/FP32/INT32/BF16 且格式为 ND。

开发者若需在自有框架的图融合 Pass 中生成该算子，可参照 test_geir_foreach_binary_op.cpp 的构图方式，并严格遵守第四节列出的约束条件。

• 人工智能
• 算子库
• 深度学习
• CANN
• Ascend