企业官网建设流程全解析

CANN ops-nn ForeachSubListInplace 算子深度解析：张量列表原地逐元素减法（x1 -= alpha·x2）的原理与 aclnn 调用实战

【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn

ForeachSubListInplace 是 CANN ops-nn 算子库foreach系列中用于"张量列表级"原地逐元素运算的算子：它接收两个张量列表x1、x2与一个标量系数alpha，对列表中每一对位置相同的 Tensor 执行x1_i = x1_i - alpha × x2_i，并将结果原地写回x1，全程不产生新输出张量。本篇文章以 foreach/foreach_sub_list_inplace/README.md 为骨架，结合算子注册定义、AICore 内核实现与单元测试源码，讲清该算子的支持范围、参数约束、两段式 aclnn 调用流程，以及底层在 NPU 上的执行原理，读完即可独立完成该算子从环境初始化到结果校验的完整 C++ 样例编写。

功能概述与计算公式

ForeachSubListInplace 解决的是"对一组结构相同、类型相同的张量批量执行减法并就地更新"的场景，与 PyTorch 的torch._foreach_sub_(...)语义一致。其核心特征是：

• 列表级（List-level）运算：输入是 Tensor 列表而非单个 Tensor，一次调用即可完成多个张量的批量减法，避免逐张量循环调用算子带来的调度开销；
• 原地更新（Inplace）：输出与第一个输入x1共享内存，计算结果直接写回x1的地址空间，不额外分配输出存储；
• 带系数（alpha）：减数侧可乘以标量系数alpha，等价于x1 = x1 - alpha * x2的融合计算。

计算公式定义如下（下标i表示列表中的第i个 Tensor，n为列表中 Tensor 数量）：

x1 = [x1_0, x1_1, ..., x1_{n-1}] x2 = [x2_0, x2_1, ..., x2_{n-1}] x1_i = x1_i - alpha * x2_i (i = 0, 1, ..., n-1)

产品支持情况

根据 foreach/foreach_sub_list_inplace/README.md，算子本身的支持情况如下：

需要注意一个细节：算子内核通路与 aclnn 接口通路的支持范围并不完全一致。算子注册定义（见下文源码分析）为 ascend950 平台配置了 AICore 内核，而 aclnnForeachSubListInplace 接口文档 中标注的接口支持情况为：Ascend 950PR/Ascend 950DT 支持，Atlas A3、Atlas A2 等平台不支持该 aclnn 接口。因此在使用时，需根据实际部署平台确认是走算子直调通路还是 aclnn 接口通路。

参数说明

算子的三个入参定义如下（信息来自 foreach/foreach_sub_list_inplace/README.md 参数表格，并与算子注册定义相互印证）：

关于参数，有几点值得展开：

• x1身兼二职：输入x1与输出x1同名，这是"原地（inplace）"语义的体现——算子在 IR 定义中不额外声明输出（详见下文op_def源码），输出直接镜像输入x1；
• alpha 的标量语义：虽然 alpha 以 Tensor 形式传入，但要求元素个数为 1，本质是一个标量系数；
• BFLOAT16 的特殊组合：当x1/x2为 BFLOAT16 时，alpha 允许使用更高精度的 FLOAT32。这一设计在算子注册定义与内核实现中均有体现（见下文内核 dispatch 分析），其原因是 BF16 的标量运算精度不足，需用 FP32 承载系数参与计算。

约束说明

• 入参x1与x2中 Tensor 的数量必须相同；
• 单个 Tensor 列表包含的 Tensor 数量不超过 256 个；
• 单个 Tensor 维度数不超过 8 维；
• x1与x2中对应位置的 Tensor 需保持数据类型、数据格式与 shape 一致；
• 同一列表内所有 Tensor 的数据类型保持一致。

aclnn 接口调用说明

按照 foreach/foreach_sub_list_inplace/README.md 的调用说明，官方推荐通过 aclnn 接口aclnnForeachSubListInplace调用该算子，完整示例见 examples/arch35/test_aclnn_foreach_sub_list_inplace.cpp，接口细节见 docs/aclnnForeachSubListInplace.md。

两段式接口函数原型

该接口遵循 CANN 标准的两段式接口设计：第一段GetWorkspaceSize完成入参校验并返回 workspace 大小与执行器，第二段真正执行计算。

aclnnStatus aclnnForeachSubListInplaceGetWorkspaceSize( aclTensorList *x1Ref, const aclTensorList *x2, const aclTensor *alpha, uint64_t *workspaceSize, aclOpExecutor **executor)

aclnnStatus aclnnForeachSubListInplace( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)

第一段接口参数说明

注意：aclnn 头文件中将被原地改写的首参数命名为x1Ref，与算子注册定义中的x1命名不同，这一细节在测试 golden 的第三方标杆适配中也有专门处理（见下文测试章节）。

返回码与错误场景

第一段接口完成入参校验，返回码遵循 aclnn 返回码 规范：

第二段接口的返回值为aclnnStatus状态码；若第一段接口已报错，不应再调用第二段接口。此外，aclnnForeachSubListInplace默认采用确定性实现。

完整调用示例

以下示例摘自 docs/aclnnForeachSubListInplace.md 并补齐了完整流程注释，编译与执行的整体流程请参考编译与运行样例：

#include <iostream> #include <vector> #include "acl/acl.h" #include "aclnnop/aclnn_foreach_sub_list_inplace.h" #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vector<int64_t>& shape) { int64_t shapeSize = 1; for (auto i : shape) { shapeSize *= i; } return shapeSize; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法，资源初始化 auto ret = aclInit(nullptr); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); ret = aclrtSetDevice(deviceId); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); ret = aclrtCreateStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); return 0; } template <typename T> int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size = GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请device侧内存 auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); // 调用aclrtMemcpy将host侧数据复制到device侧内存上 ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); // 计算连续tensor的strides std::vector<int64_t> strides(shape.size(), 1); for (int64_t i = shape.size() - 2; i >= 0; i--) { strides[i] = shape[i + 1] * strides[i + 1]; } // 调用aclCreateTensor接口创建aclTensor *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; } int main() { // 1. （固定写法）device/stream初始化，参考acl API手册 // 根据自己的实际device填写deviceId int32_t deviceId = 0; aclrtStream stream; auto ret = Init(deviceId, &stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); // 2. 构造输入与输出，需要根据API的接口自定义构造 std::vector<int64_t> selfShape1 = {2, 3}; std::vector<int64_t> selfShape2 = {1, 3}; std::vector<int64_t> otherShape1 = {2, 3}; std::vector<int64_t> otherShape2 = {1, 3}; std::vector<int64_t> alphaShape = {1}; void* input1DeviceAddr = nullptr; void* input2DeviceAddr = nullptr; void* other1DeviceAddr = nullptr; void* other2DeviceAddr = nullptr; void* alphaDeviceAddr = nullptr; aclTensor* input1 = nullptr; aclTensor* input2 = nullptr; aclTensor* other1 = nullptr; aclTensor* other2 = nullptr; aclTensor* alpha = nullptr; std::vector<float> input1HostData = {3, 5, 7, 4, 5, 9}; std::vector<float> input2HostData = {5, 4, 1}; std::vector<float> other1HostData = {1, 2, 3, 4, 5, 6}; std::vector<float> other2HostData = {7, 8, 9}; std::vector<float> alphaValueHostData = {1.2f}; // 创建input1 aclTensor ret = CreateAclTensor(input1HostData, selfShape1, &input1DeviceAddr, aclDataType::ACL_FLOAT, &input1); CHECK_RET(ret == ACL_SUCCESS, return ret); // 创建input2 aclTensor ret = CreateAclTensor(input2HostData, selfShape2, &input2DeviceAddr, aclDataType::ACL_FLOAT, &input2); CHECK_RET(ret == ACL_SUCCESS, return ret); // 创建other1 aclTensor ret = CreateAclTensor(other1HostData, otherShape1, &other1DeviceAddr, aclDataType::ACL_FLOAT, &other1); CHECK_RET(ret == ACL_SUCCESS, return ret); // 创建other2 aclTensor ret = CreateAclTensor(other2HostData, otherShape2, &other2DeviceAddr, aclDataType::ACL_FLOAT, &other2); CHECK_RET(ret == ACL_SUCCESS, return ret); // 创建alpha aclTensor ret = CreateAclTensor(alphaValueHostData, alphaShape, &alphaDeviceAddr, aclDataType::ACL_FLOAT, &alpha); CHECK_RET(ret == ACL_SUCCESS, return ret); std::vector<aclTensor*> tempInput1{input1, input2}; aclTensorList* tensorListInput1 = aclCreateTensorList(tempInput1.data(), tempInput1.size()); std::vector<aclTensor*> tempInput2{other1, other2}; aclTensorList* tensorListInput2 = aclCreateTensorList(tempInput2.data(), tempInput2.size()); // 3. 调用CANN算子库API，需要修改为具体的API名称 uint64_t workspaceSize = 0; aclOpExecutor* executor; // 调用aclnnForeachSubListInplace第一段接口 ret = aclnnForeachSubListInplaceGetWorkspaceSize(tensorListInput1, tensorListInput2, alpha, &workspaceSize, &executor); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnForeachSubListInplaceGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* workspaceAddr = nullptr; if (workspaceSize > 0) { ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); } // 调用aclnnForeachSubListInplace第二段接口 ret = aclnnForeachSubListInplace(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnForeachSubListInplace failed. ERROR: %d\n", ret); return ret); // 4. （固定写法）同步等待任务执行结束 ret = aclrtSynchronizeStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); // 5. 获取输出的值，将device侧内存上的结果复制至host侧，需要根据具体API的接口定义修改 auto size = GetShapeSize(selfShape1); std::vector<float> self1Data(size, 0); ret = aclrtMemcpy(self1Data.data(), self1Data.size() * sizeof(self1Data[0]), input1DeviceAddr, size * sizeof(self1Data[0]), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); for (int64_t i = 0; i < size; i++) { LOG_PRINT("out1 result[%ld] is: %f\n", i, self1Data[i]); } size = GetShapeSize(selfShape2); std::vector<float> self2Data(size, 0); ret = aclrtMemcpy(self2Data.data(), self2Data.size() * sizeof(self2Data[0]), input2DeviceAddr, size * sizeof(self2Data[0]), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); for (int64_t i = 0; i < size; i++) { LOG_PRINT("out2 result[%ld] is: %f\n", i, self2Data[i]); } // 6. 释放aclTensor，需要根据具体API的接口定义修改 aclDestroyTensorList(tensorListInput1); aclDestroyTensorList(tensorListInput2); aclDestroyTensor(alpha); // 7.释放device资源，需要根据具体API的接口定义修改 aclrtFree(input1DeviceAddr); aclrtFree(input2DeviceAddr); aclrtFree(other1DeviceAddr); aclrtFree(other2DeviceAddr); aclrtFree(alphaDeviceAddr); if (workspaceSize > 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }

示例中x1列表包含两个 Tensor（shape 分别为{2,3}与{1,3}），x2列表包含两个 Tensor，alpha = 1.2。由于是原地运算，最终结果直接读取input1DeviceAddr、input2DeviceAddr即可：第一个 Tensor 结果为3-1.2×1=1.8、5-1.2×2=2.6、7-1.2×3=3.4、4-1.2×4=-0.8、5-1.2×5=-1、9-1.2×6=1.8，第二个 Tensor 结果为5-1.2×7=-3.4、4-1.2×8=-5.6、1-1.2×9=-9.8。

源码级实现原理

算子注册定义：inplace 语义的 IR 表达

在 op_host/foreach_sub_list_inplace_def.cpp 中，算子通过OpDef注册，其注释直接点明设计要点："Inplace: x1 = x1 - alpha * x2, x1 serves as both input and output (no Output declared)"——输入输出共享x1名称，输出x1的ParamType(DYNAMIC)与输入x1一致，实现"同名镜像"的原地语义：

class ForeachSubListInplace : public OpDef { public: explicit ForeachSubListInplace(const char* name) : OpDef(name) { this->Input("x1") .ParamType(DYNAMIC) .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_BF16}) .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .AutoContiguous(); this->Input("x2") .ParamType(DYNAMIC) .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_BF16}) .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .AutoContiguous(); this->Input("alpha") .ParamType(REQUIRED) .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_FLOAT}) .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .AutoContiguous(); this->Output("x1") .ParamType(DYNAMIC) .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_BF16}) .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) .AutoContiguous(); OpAICoreConfig regbaseCfg; regbaseCfg.DynamicCompileStaticFlag(true).DynamicRankSupportFlag(true).DynamicShapeSupportFlag(true); this->AICore().AddConfig("ascend950", regbaseCfg); } }; OP_ADD(ForeachSubListInplace);

从注册定义可以确认几点实现事实：x1、x2均声明为 DYNAMIC 输入（Tensor 列表数量运行时确定）；alpha为 REQUIRED 标量输入，其DataType列表{FLOAT16, FLOAT, INT32, FLOAT}中第 4 项重复为FLOAT，正对应"BFLOAT16 输入时 alpha 用 FLOAT32"的规则（x1为 BF16 时按位序对应 FP32 的 alpha）；支持动态 shape、动态 rank 与动态编译，数据格式统一为 ND。

IR 原型定义

op_graph/foreach_sub_list_inplace_proto.h 中通过REG_OP宏给出与op_def一致的算子原型，头部注释同样给出公式 "y(=x1) = x1 - alpha * x2, element-wise over two tensor lists, written back in place to x1"：

REG_OP(ForeachSubListInplace) .DYNAMIC_INPUT(x1, TensorType({DT_FLOAT, DT_FLOAT16, DT_INT32, DT_BF16})) .DYNAMIC_INPUT(x2, TensorType({DT_FLOAT, DT_FLOAT16, DT_INT32, DT_BF16})) .INPUT(alpha, TensorType({DT_FLOAT, DT_FLOAT16, DT_INT32})) .DYNAMIC_OUTPUT(x1, TensorType({DT_FLOAT, DT_FLOAT16, DT_INT32, DT_BF16})) .OP_END_FACTORY_REG(ForeachSubListInplace)

AICore 内核实现：类型分发与精度策略

内核代码位于 op_kernel/arch35/foreach_sub_list_inplace.cpp，入口函数根据 tiling key 将不同数据类型分发给对应模板实例：

extern "C" __global__ __aicore__ void foreach_sub_list_inplace(GM_ADDR x1, GM_ADDR x2, GM_ADDR alpha, GM_ADDR x1Ref, GM_ADDR workspace, GM_ADDR tiling) { TPipe pipeOp; if (TILING_KEY_IS(FOREACH_TILING_KEY_HALF)) { ForeachSubListInplaceImpl<half, DTYPE_ALPHA>(x1, x2, alpha, x1Ref, workspace, tiling, &pipeOp); } else if (TILING_KEY_IS(FOREACH_TILING_KEY_FLOAT)) { ForeachSubListInplaceImpl<float, float>(x1, x2, alpha, x1Ref, workspace, tiling, &pipeOp); } else if (TILING_KEY_IS(FOREACH_TILING_KEY_INT)) { ForeachSubListInplaceImpl<int, int>(x1, x2, alpha, x1Ref, workspace, tiling, &pipeOp); } else if (TILING_KEY_IS(FOREACH_TILING_KEY_BF16)) { ForeachSubListInplaceImpl<bfloat16_t, float>(x1, x2, alpha, x1Ref, workspace, tiling, &pipeOp); } pipeOp.Destroy(); }

从内核的模板实例化可以清晰看到四条计算路径与精度策略：

• FLOAT16（half）：元素类型为half，alpha 类型为DTYPE_ALPHA（即 half），半精度输入、半精度系数直接计算；
• FLOAT32（float）：<float, float>，原生 FP32 计算；
• INT32（int）：<int, int>，整型输入、整型系数，执行整数减法（测试 golden 中特别强调整型路径遵循补码回绕语义，不做饱和处理）；
• BFLOAT16（bfloat16_t）：<bfloat16_t, float>——元素为 BF16 但 alpha 提升为 FP32 参与计算，这正是参数表中"x1为 BFLOAT16 时 alpha 支持 FLOAT32"的内核侧体现。

内核真正的逐元素运算逻辑由 op_kernel/arch35/foreach_sub_list_inplace_regbase.h 提供，它复用 foreach 系列的公共基类ForeachBinaryAlphaCastInplaceRegbase（定义于 foreach/foreach_utils 目录），仅需注入算子特有的二元运算即可：

template <typename T, typename ScalarT, typename Tiling> class ForeachSubListInplaceRegbase : public ForeachBinaryAlphaCastInplaceRegbase<T, ScalarT, Tiling, ForeachSubListInplaceRegbase<T, ScalarT, Tiling>> { public: template <typename U> __aicore__ inline void ApplyBinaryOp(LocalTensor<U> dst, LocalTensor<U> a, LocalTensor<U> b, int64_t dataCount) { Sub(dst, a, b, dataCount); } };

基类负责 tiling 数据解析、Init、Process以及"先乘 alpha 再做二元运算"的公共流水，子类只实现Sub(dst, a, b, dataCount)一步减法。注释还披露了一个平台细节：dav-3510 上不支持 BF16 标量 cast，因此 BF16 路径必须将标量按 FP32 处理。

Host 侧配置：编译产物与数据类型对应

op_host/config/ascend950/foreach_sub_list_inplace_binary.json 列出了该算子在 ascend950 平台预编译二进制与数据类型的对应关系：共 4 组 binary，分别对应 float16、float32、int32、bfloat16 四种x1/x2组合；其中 bfloat16 组合的 alpha 声明为float32，与注册定义和内核 dispatch 完全一致。所有输入的shape均为[-2]（表示动态 shape），format_match_mode为FormatAgnostic。

测试验证：infershape 单测与 golden 精度对齐

算子目录下有两类测试从不同层面保障正确性：

Host 侧 infershape 单测tests/ut/op_host/test_foreach_sub_list_inplace_infershape.cpp 验证了两件事：一是 shape 推导，输入列表包含 6 个 shape 分别为[8]~[13]的 Tensor（IR 实例数{3, 3, 1}对应 x1 3 个、x2 3 个、alpha 1 个），输出 shape 必须被真正推导为与 x1 相同的[8]、[9]、[10]，验证"就地参数由同名 IR 输出镜像（ref 配对）"的语义；二是数据类型推导，输入为 FLOAT16 时输出类型推导成功返回GRAPH_SUCCESS。

精度 goldentests/assets/golden.py 使用 PyTorch 的torch._foreach_mul+torch._foreach_sub两步拼接复现算子计算，而非单步_foreach_sub(alpha=)的 FMA 形式。这一选择有明确的精度动机：NPU 内核是"先 Muls 再 Sub"的两次舍入，若 golden 用 FMA 单次舍入，误差分布与内核不一致会导致比对假红。golden 同时覆盖了整型补码回绕语义（int64 中间量 + 窄化回 int32）与 BF16 的 FP32 加宽计算口径，并将torch作为三方标杆（third_party）用于 cross_check 比对。

总结

ForeachSubListInplace 是 CANN ops-nn 中面向张量列表的原地减法算子，通过x1_i = x1_i - alpha * x2_i一次调用完成批量张量的带系数逐元素减法。本文梳理了其产品支持范围（算子通路支持 Ascend 950 及 Atlas A2/A3 系列，aclnn 接口通路当前仅支持 Ascend 950 系列）、参数约束（列表 Tensor 数量一致且不超过 256 个、维度不超过 8 维、alpha 为单元素标量）、两段式 aclnn 调用流程与完整示例代码，并从注册定义、IR 原型、AICore 内核 dispatch、binary 配置到 golden 测试五个层面剖析了其实现原理。对于需要在 NPU 上对多组张量执行批量减法并节省输出内存的开发者而言，该算子提供了接口简洁、语义明确、内存高效的解决方案。

【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn