CANN ops-nn ForeachSubListInplace 算子深度解析张量列表原地逐元素减法x1 - alpha·x2的原理与 aclnn 调用实战【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nnForeachSubListInplace 是 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个 Tensorn为列表中 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算子本身的支持情况如下产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品×Atlas 训练系列产品×Kirin X90 处理器系列产品×Kirin 9030 处理器系列产品×需要注意一个细节算子内核通路与 aclnn 接口通路的支持范围并不完全一致。算子注册定义见下文源码分析为 ascend950 平台配置了 AICore 内核而 aclnnForeachSubListInplace 接口文档 中标注的接口支持情况为Ascend 950PR/Ascend 950DT 支持Atlas A3、Atlas A2 等平台不支持该 aclnn 接口。因此在使用时需根据实际部署平台确认是走算子直调通路还是 aclnn 接口通路。参数说明算子的三个入参定义如下信息来自 foreach/foreach_sub_list_inplace/README.md 参数表格并与算子注册定义相互印证参数名输入/输出描述数据类型数据格式x1输入支持空 Tensor。表示相减运算的第一个输入张量列表对应公式中的x1。该参数中所有 Tensor 的数据类型保持一致每个 Tensor 的维度数不超过 8 维。FLOAT32、FLOAT16、INT32、BFLOAT16NDx1输出支持空 Tensor。表示计算结果张量列表与输入x1共享内存计算结果原地写回。该参数中每个 Tensor 的 shape 和数据类型均与输入x1中对应位置的 Tensor 一致。FLOAT32、FLOAT16、INT32、BFLOAT16NDx2输入支持空 Tensor。表示相减运算的第二个输入张量列表对应公式中的x2。数据类型、数据格式和 shape 与入参x1一致。该参数中所有 Tensor 的数据类型保持一致。FLOAT32、FLOAT16、INT32、BFLOAT16NDalpha输入不支持空 Tensor。表示相减运算中第二个输入的系数对应公式中的alpha。元素个数为 1。数据类型与入参x1具有一定对应关系当x1为 FLOAT32、FLOAT16、INT32 时alpha 与x1类型一致当x1为 BFLOAT16 时alpha 支持 FLOAT32。FLOAT32、FLOAT16、INT32ND关于参数有几点值得展开x1身兼二职输入x1与输出x1同名这是原地inplace语义的体现——算子在 IR 定义中不额外声明输出详见下文op_def源码输出直接镜像输入x1alpha 的标量语义虽然 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)第一段接口参数说明参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续 Tensorx1RefaclTensorList*输入/输出被减数输入和输出张量列表对应公式中的x1支持空 tensor所有 Tensor 数据类型保持一致FLOAT32、FLOAT16、BFLOAT16、INT32ND0-8√x2aclTensorList*输入减数张量列表对应公式中的x2支持空 tensor所有 Tensor 数据类型保持一致数据类型、数据格式和 shape 与x1Ref一致FLOAT32、FLOAT16、BFLOAT16、INT32ND0-8√alphaaclTensor*输入减数系数对应公式中的alpha不支持空 tensor元素个数为 1x1Ref为 FLOAT32/FLOAT16/INT32 时类型一致x1Ref为 BFLOAT16 时支持 FLOAT32FLOAT32、FLOAT16、INT32ND0-8√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小-----executoraclOpExecutor**输出返回 op 执行器包含算子计算流程-----注意aclnn 头文件中将被原地改写的首参数命名为x1Ref与算子注册定义中的x1命名不同这一细节在测试 golden 的第三方标杆适配中也有专门处理见下文测试章节。返回码与错误场景第一段接口完成入参校验返回码遵循 aclnn 返回码 规范返回码错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 x1Ref、x2、alpha 是空指针ACLNN_ERR_PARAM_INVALID161002x1Ref、x2、alpha 的数据类型不在支持范围之内ACLNN_ERR_PARAM_INVALID161002x1Ref、x2 的数据类型不一致ACLNN_ERR_PARAM_INVALID161002x1Ref、x2 中存在空指针 TensorACLNN_ERR_INNER_TILING_ERROR561002x1Ref、x2 的 shape 不满足约束ACLNN_ERR_INNER_TILING_ERROR561002x1Ref、x2 中的 Tensor 的数据类型不一致ACLNN_ERR_INNER_TILING_ERROR561002x1Ref、x2 中的 Tensor 维度超过 8 维ACLNN_ERR_INNER_TILING_ERROR561002alpha 元素个数不为 1ACLNN_ERR_INNER_TILING_ERROR561002x1Ref、x2 的 Tensor 数量不一致第二段接口的返回值为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::vectorint64_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::vectorT hostData, const std::vectorint64_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::vectorint64_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::vectorint64_t selfShape1 {2, 3}; std::vectorint64_t selfShape2 {1, 3}; std::vectorint64_t otherShape1 {2, 3}; std::vectorint64_t otherShape2 {1, 3}; std::vectorint64_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::vectorfloat input1HostData {3, 5, 7, 4, 5, 9}; std::vectorfloat input2HostData {5, 4, 1}; std::vectorfloat other1HostData {1, 2, 3, 4, 5, 6}; std::vectorfloat other2HostData {7, 8, 9}; std::vectorfloat 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::vectoraclTensor* tempInput1{input1, input2}; aclTensorList* tensorListInput1 aclCreateTensorList(tempInput1.data(), tempInput1.size()); std::vectoraclTensor* 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::vectorfloat 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::vectorfloat 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列表包含两个 Tensorshape 分别为{2,3}与{1,3}x2列表包含两个 Tensoralpha 1.2。由于是原地运算最终结果直接读取input1DeviceAddr、input2DeviceAddr即可第一个 Tensor 结果为3-1.2×11.8、5-1.2×22.6、7-1.2×33.4、4-1.2×4-0.8、5-1.2×5-1、9-1.2×61.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 x1REG_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)) { ForeachSubListInplaceImplhalf, DTYPE_ALPHA(x1, x2, alpha, x1Ref, workspace, tiling, pipeOp); } else if (TILING_KEY_IS(FOREACH_TILING_KEY_FLOAT)) { ForeachSubListInplaceImplfloat, float(x1, x2, alpha, x1Ref, workspace, tiling, pipeOp); } else if (TILING_KEY_IS(FOREACH_TILING_KEY_INT)) { ForeachSubListInplaceImplint, int(x1, x2, alpha, x1Ref, workspace, tiling, pipeOp); } else if (TILING_KEY_IS(FOREACH_TILING_KEY_BF16)) { ForeachSubListInplaceImplbfloat16_t, float(x1, x2, alpha, x1Ref, workspace, tiling, pipeOp); } pipeOp.Destroy(); }从内核的模板实例化可以清晰看到四条计算路径与精度策略FLOAT16half元素类型为halfalpha 类型为DTYPE_ALPHA即 half半精度输入、半精度系数直接计算FLOAT32floatfloat, float原生 FP32 计算INT32intint, int整型输入、整型系数执行整数减法测试 golden 中特别强调整型路径遵循补码回绕语义不做饱和处理BFLOAT16bfloat16_tbfloat16_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 ForeachBinaryAlphaCastInplaceRegbaseT, ScalarT, Tiling, ForeachSubListInplaceRegbaseT, ScalarT, Tiling { public: template typename U __aicore__ inline void ApplyBinaryOp(LocalTensorU dst, LocalTensorU a, LocalTensorU 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]表示动态 shapeformat_match_mode为FormatAgnostic。测试验证infershape 单测与 golden 精度对齐算子目录下有两类测试从不同层面保障正确性Host 侧 infershape 单测tests/ut/op_host/test_foreach_sub_list_inplace_infershape.cpp 验证了两件事一是 shape 推导输入列表包含 6 个 shape 分别为[8]~[13]的 TensorIR 实例数{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_multorch._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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考