人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载本篇技术指南以 CANN 神经网络算子库 ops-nn 中的 ForeachErfc 算子为对象系统讲解其功能定义对输入张量列表逐元素执行互补误差函数 erfc 运算、产品支持范围、参数约束、两种调用方式aclnn 接口与图模式构图并结合仓库中的算子 IR、Host 侧 Tiling 实现与 Kernel 源码深入剖析其底层实现原理。读完本文读者可以掌握在 Ascend NPU 与 Kirin 处理器上通过aclnnForeachErfc两段式接口对一组 Tensor 批量执行 erfc 计算的完整实战方法。产品支持情况ForeachErfc 算子在不同硬件产品上的支持情况如下表所示数据来源于 foreach/foreach_erfc/README.md产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品×Atlas 训练系列产品×Kirin X90 处理器系列产品√Kirin 9030 处理器系列产品√从源码层面看这一支持范围与算子在 Host 侧算子定义 foreach_erfc_def.cpp 中注册的 AICore 配置一一对应算子为ascend950、ascend910_93即 Atlas A3 系列、ascend910b即 Atlas A2 系列、kirinx90与kirin9030分别注册了对应的 AICore 配置而未支持的产品如 Atlas 训练系列、推理系列、Atlas 200I/500 A2则未注册对应配置。同时不同产品在数据类型支持上存在差异需特别留意Kirin X90/Kirin 9030 处理器系列产品不支持 BFLOAT16。这一限制同样体现在算子定义代码中kirinx90/kirin9030使用的 Kirin 专属配置仅注册了DT_FLOAT16、DT_FLOAT两种数据类型见 foreach_erfc_def.cpp而其他产品的默认配置支持DT_FLOAT16、DT_FLOAT、DT_BF16三种类型。功能说明ForeachErfc 是 foreach 系列算子族的一员其核心功能是按元素对输入张量列表中的每个 Tensor 执行互补误差函数complementary error function运算。算子功能按元素执行从 x 到无穷大积分的互补误差函数运算。计算公式$$ x [{x_0}, {x_1}, ... {x_{n-1}}]\ y [{y_0}, {y_1}, ... {y_{n-1}}] $$$$ y_i {\rm erfc}(x_i) 1 - {\rm erf}(x_i) 1 - \frac{2}{\sqrt{\pi}} \int_{0}^{x_i} e^{-t^2} dt \frac{2}{\sqrt{\pi}} \int_{x_i}^{\infty} e^{-t^2} dt (i0,1,...n-1) $$其中x与y均为 Tensor 列表TensorList列表中的第i个输入张量x_i逐元素计算出第i个输出张量y_i。互补误差函数erfc(x) 1 - erf(x)在概率论与统计学中常用于高斯分布尾部概率的计算如正态分布的右尾积分与 erf 误差函数互为补余关系。该算子的计算由 Kernel 侧调用 AscendC 数学库中的Erfc模板完成。在 foreach_erfc.cpp 中可以看到Kernel 通过ForeachTrianglesrcT, dstT, ErfcT, false模板对输入张量列表进行统一的遍历与搬运并按数据类型TilingKey分发到不同的实例化路径TILING_KEY_IS(1)ForeachTrianglehalf, half, Erfchalf, false对应 FLOAT16 输入、FLOAT16 输出TILING_KEY_IS(2)ForeachTrianglefloat, float, Erfcfloat, false对应 FLOAT32 输入、FLOAT32 输出TILING_KEY_IS(4)ForeachTrianglebfloat16_t, float, Erfcfloat, false对应 BFLOAT16 输入、FLOAT32 中间精度计算后输出 BFLOAT16该分支在__NPU_ARCH__为 3003/3113 的硬件上不参与编译。在 Ascend 950arch35平台上算子另有一套基于 SIMT 架构的实现 arch35/foreach_erfc.cpp通过NsForeachErfc::Processhalf/float/bfloat16_t分别处理三种数据类型。值得一提的是BFLOAT16 输入在计算时以 FLOAT32 精度参与中间运算再回写为 BFLOAT16这有助于缓解低精度数据类型带来的数值误差。参数说明ForeachErfc 算子的输入输出参数如下内容继承自 foreach/foreach_erfc/README.md 并补充了 aclnn 接口层面的细节参见 docs/aclnnForeachErfc.md参数名输入/输出/属性描述数据类型数据格式x输入表示进行误差函数运算的输入张量列表对应公式中的x。该参数中所有 Tensor 的数据类型保持一致。FLOAT32、FLOAT16、BFLOAT16NDy输出表示进行误差函数运算的输出张量列表对应公式中的y。该参数中所有 Tensor 的数据类型保持一致。数据类型和数据格式与入参x的数据类型和数据格式一致shape size 大于等于入参x的 shape size。FLOAT32、FLOAT16、BFLOAT16ND补充说明在 aclnn 接口aclnnForeachErfcGetWorkspaceSize层面输入x与输出out均为aclTensorList*类型支持空 Tensor 列表x支持非连续 Tensor而out不支持非连续 Tensorshape 维度范围为 0~8 维。从算子 IR 定义 foreach_erfc_proto.h 可以看到x声明为DYNAMIC_INPUT、y声明为DYNAMIC_OUTPUT即输入/输出 Tensor 数量在运行时动态确定这正是 foreach 系列算子支持任意数量 Tensor 列表的协议基础。在算子定义 foreach_erfc_def.cpp 中输入输出均配置为AutoContiguous()并开启了DynamicCompileStaticFlag(true)、DynamicRankSupportFlag(true)、DynamicShapeSupportFlag(true)等动态特性表明算子支持动态 shape、动态 rank0~8 维与动态编译。底层二进制配置可参见各平台目录下的 foreach_erfc_binary.json其中按 float16、float32、bfloat16 三种数据类型分别注册了对应的 kernel binarybin_filenameparamType均为dynamic、format均为 ND与上表描述一致。约束说明输出不支持非连续 Tensor即输出y列表中的每个 Tensor 必须是内存连续contiguous的。相比之下输入x在 aclnn 接口层面支持非连续 Tensor计算框架会在必要时自动搬运为连续内存。数据类型一致性x列表内部、y列表内部以及x与y之间的数据类型必须保持一致。shape 关系输出y的 shape size 需大于等于输入x的 shape size。Kirin 平台数据类型限制Kirin X90/Kirin 9030 不支持 BFLOAT16仅支持 FLOAT16 与 FLOAT32。上述约束在 aclnn 接口第一段入参校验中均会触发对应错误码详见下文返回码小节。调用说明ForeachErfc 算子支持两种调用方式调用方式样例代码说明aclnn 接口test_aclnn_foreach_erfc通过aclnnForeachErfc接口方式调用 ForeachErfc 算子。图模式-通过算子 IR 构图方式调用 ForeachErfc 算子。aclnn 两段式接口aclnnAscend CANN Lite 神经网络接口采用两段式调用必须先调用aclnnForeachErfcGetWorkspaceSize接口获取计算所需 workspace 大小以及包含了算子计算流程的执行器executor再调用aclnnForeachErfc接口执行计算。函数原型如下aclnnStatus aclnnForeachErfcGetWorkspaceSize( const aclTensorList *x, const aclTensorList *out, uint64_t *workspaceSize, aclOpExecutor **executor)aclnnStatus aclnnForeachErfc( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)第一段接口aclnnForeachErfcGetWorkspaceSize参数说明参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续 TensorxaclTensorList*输入表示按元素返回从 x 到无穷大积分的互补误差函数运算的输入张量列表对应公式中的x。支持空 Tensor该参数中所有 Tensor 的数据类型保持一致。BFLOAT16、FLOAT16、FLOAT32ND0-8√outaclTensorList*输出表示按元素返回从 x 到无穷大积分的互补误差函数运算的输出张量列表对应公式中的y。支持空 Tensor该参数中所有 Tensor 的数据类型保持一致数据类型和数据格式与入参x一致shape size 大于等于入参x的 shape size。BFLOAT16、FLOAT16、FLOAT32ND0-8×workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小。-----executoraclOpExecutor**输出返回 op 执行器包含了算子计算流程。-----返回码aclnnStatus为返回状态码。第一段接口完成入参校验出现以下场景时报错返回码错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 x 或 out 是空指针。ACLNN_ERR_PARAM_INVALID161002x 或 out 的数据类型不在支持的范围之内。ACLNN_ERR_PARAM_INVALID161002x 和 out 的数据类型不一致。ACLNN_ERR_INNER_TILING_ERROR561002x 与 out 的 shape 不满足约束。ACLNN_ERR_INNER_TILING_ERROR561002x 或 out 中的 Tensor 的数据类型不一致。ACLNN_ERR_INNER_TILING_ERROR561002x 或 out 中的 Tensor 维度超过 8 维。第二段接口aclnnForeachErfc参数说明参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址。workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口aclnnForeachErfcGetWorkspaceSize获取。executor输入op 执行器包含了算子计算流程。stream输入指定执行任务的 Stream。返回值aclnnStatus状态码。确定性计算aclnnForeachErfc默认确定性实现即相同输入在多次运行下产生确定一致的输出结果。完整调用示例代码以下示例来自仓库 examples/test_aclnn_foreach_erfc.cpp演示了对包含两个 Tensorshape 分别为{2, 3}与{1, 3}的输入列表批量执行 erfc 的完整流程。代码整体分为设备/Stream 初始化 → 构造输入输出 → 两段式调用 → 同步等待 → 取回结果 → 资源释放六个步骤#include iostream #include vector #include acl/acl.h #include aclnnop/aclnn_foreach_erfc.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手册 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. 构造输入与输出 std::vectorint64_t selfShape1 {2, 3}; std::vectorint64_t selfShape2 {1, 3}; std::vectorint64_t outShape1 {2, 3}; std::vectorint64_t outShape2 {1, 3}; void* input1DeviceAddr nullptr; void* input2DeviceAddr nullptr; void* out1DeviceAddr nullptr; void* out2DeviceAddr nullptr; aclTensor* input1 nullptr; aclTensor* input2 nullptr; aclTensor* out1 nullptr; aclTensor* out2 nullptr; std::vectorfloat input1HostData {1, 2, 3, 4, 5, 6}; std::vectorfloat input2HostData {7, 8, 9}; std::vectorfloat out1HostData(6, 0); std::vectorfloat out2HostData(3, 0); // 创建input1/input2/out1/out2 aclTensor ret CreateAclTensor(input1HostData, selfShape1, input1DeviceAddr, aclDataType::ACL_FLOAT, input1); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(input2HostData, selfShape2, input2DeviceAddr, aclDataType::ACL_FLOAT, input2); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(out1HostData, outShape1, out1DeviceAddr, aclDataType::ACL_FLOAT, out1); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(out2HostData, outShape2, out2DeviceAddr, aclDataType::ACL_FLOAT, out2); CHECK_RET(ret ACL_SUCCESS, return ret); // 组装TensorList std::vectoraclTensor* tempInput{input1, input2}; aclTensorList* tensorListInput aclCreateTensorList(tempInput.data(), tempInput.size()); std::vectoraclTensor* tempOutput{out1, out2}; aclTensorList* tensorListOutput aclCreateTensorList(tempOutput.data(), tempOutput.size()); // 3. 两段式调用CANN算子库API uint64_t workspaceSize 0; aclOpExecutor* executor; // 第一段接口计算workspace大小并获取执行器 ret aclnnForeachErfcGetWorkspaceSize(tensorListInput, tensorListOutput, workspaceSize, executor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnForeachErfcGetWorkspaceSize 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); } // 第二段接口执行计算 ret aclnnForeachErfc(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnForeachErfc 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侧并打印 auto size GetShapeSize(outShape1); std::vectorfloat out1Data(size, 0); ret aclrtMemcpy(out1Data.data(), out1Data.size() * sizeof(out1Data[0]), out1DeviceAddr, size * sizeof(out1Data[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, out1Data[i]); } size GetShapeSize(outShape2); std::vectorfloat out2Data(size, 0); ret aclrtMemcpy(out2Data.data(), out2Data.size() * sizeof(out2Data[0]), out2DeviceAddr, size * sizeof(out2Data[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, out2Data[i]); } // 6. 释放aclTensor与device资源 aclDestroyTensorList(tensorListInput); aclDestroyTensorList(tensorListOutput); aclrtFree(input1DeviceAddr); aclrtFree(input2DeviceAddr); aclrtFree(out1DeviceAddr); aclrtFree(out2DeviceAddr); if (workspaceSize 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }示例中的几个关键点说明TensorList 的组装使用aclCreateTensorList将多个aclTensor*打包为aclTensorList*这是 foreach 系列算子区别于单 Tensor 算子的核心入参形态。列表内各 Tensor 的 shape 可以不同示例中为{2, 3}与{1, 3}只要数据类型一致即可。workspace 申请仅当第一段接口返回的workspaceSize 0时才需要aclrtMalloc申请内存返回 0 时可跳过。数据类型的切换示例中使用aclDataType::ACL_FLOAT与std::vectorfloat若改用 FLOAT16则对应ACL_HALF与std::vectorhalf或aclfp16即可。图模式调用图模式通过算子 IR 构图方式调用。ForeachErfc 的算子原型定义位于 op_graph/foreach_erfc_proto.h其核心内容如下REG_OP(ForeachErfc) .DYNAMIC_INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16, DT_BF16})) .DYNAMIC_OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16, DT_BF16})) .OP_END_FACTORY_REG(ForeachErfc)即算子名为ForeachErfc输入x与输出y均为动态个数的 Tensor 列表数据类型支持DT_FLOAT、DT_FLOAT16、DT_BF16。在图模式下开发者通过 GEGraph Engine构图接口创建ForeachErfc算子节点将若干 Tensor 挂接为x的动态输入图编译执行时即会调度到该算子对应的 Kernel 完成 erfc 批量计算。从源码理解算子实现链路Host 侧 Tiling切分策略与核心参数Tiling数据切分是算子能否高效利用多核的关键。以 arch35Ascend 950平台为例Tiling 实现在 op_host/arch35/foreach_erfc_tiling_arch35.cpp平台信息获取通过GetPlatformInfoFallback读取 AIV 核数coreNum与 UB 内存大小ubSizeUB 大小需大于 128KBDCACHE_SIZE的保留缓存实际可用 UB 为ubSize - DCACHE_SIZE。Tensor 列表展平遍历输入列表中的每个 Tensor累加其元素个数得到totalElements并将每个 Tensor 的累计偏移cumulativeOffsets[i]写入 TilingData。这一设计将Tensor 列表逻辑上展平为一段连续数据Kernel 侧根据累计偏移即可定位每个子 Tensor 的起始位置。多核切分按SINGLE_CORE_MIN_ELEMENTS 1024的粒度估算所需核数即needCoreNum min(coreNum, ceil(totalElements / 1024))再计算每个核处理的元素数perCoreElements ceil(totalElements / needCoreNum)并通过context-SetBlockDim(needCoreNum)设置实际启动的核数。TilingKey 与 Workspace根据输入首 Tensor 的数据类型设置 TilingKeyFLOAT16→0、FLOAT32→1、BFLOAT16→2用于 Kernel 侧选择对应的实例化路径同时将 workspace 大小置 0即该算子在 arch35 平台上无需额外 workspace。单测 tests/ut/op_host/arch22/test_foreach_erfc_tiling.cpp 与 tests/ut/op_host/arch35/test_foreach_erfc_tiling.cpp 对上述 Tiling 逻辑进行了验证测试通过gert::OpImplRegistry获取注册的 Tiling 函数与解析函数并构造平台信息UB_SIZE、CORE_NUM 等后运行 Tiling 并检查输出结果。Kernel 侧计算ForeachTriangle 与 ErfcKernel 入口在 op_kernel/foreach_erfc.cpp对外导出extern C __global__ __aicore__ void foreach_erfc(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)。其内部逻辑为通过GET_TILING_DATA解析 Host 侧下发的 TilingData根据 TilingKey 选择数据类型对应的ForeachTrianglesrcT, dstT, ErfcT, false模板实例op.Init(...)完成地址与 TilingData 的初始化绑定op.Process()按累计偏移遍历各子 Tensor逐元素调用 AscendC 数学库lib/math/kernel_operator_erf_intf.h中的Erfc完成互补误差函数计算。ForeachTriangle是 foreach 系列算子共用的张量列表遍历框架定义于foreach_utils它负责将动态数量的 Tensor 列表按 TilingData 中的偏移切分到各核并驱动搬运与计算Erfc则封装了具体的数学计算。这种通用遍历框架 算子专属计算函数的组合是 foreach 算子族高效复用的关键设计。测试验证仓库为 ForeachErfc 提供了完整的单测UT与系统测试ST覆盖Kernel 单测tests/ut/op_kernel/test_foreach_erfc.cpp 通过 tikit 的 CPU 仿真模式ICPU_RUN_KF直接运行 Kernel并使用 tests/ut/op_kernel/erfc_data/gen_data.py 生成 fp16/bf16 等类型的测试数据计算完成后由 compare_data.py 与参考结果对比校验。测试用例覆盖了 FLOAT16、BFLOAT16 等数据类型shape 组合如{{128, 64}, {16, 128}, {32, 128}}的三 Tensor 输入列表。ST 测试tests/st/aclnnForeachErfc/ 下提供了基于 ATK 框架的用例配置文件atk_aclnnForeachErfc.json与执行脚本executor_aclnnForeachErfc.py用于在真实 NPU 环境上验证 aclnn 接口的行为。小结ForeachErfc 是 CANN ops-nn 中用于对 Tensor 列表逐元素执行互补误差函数erfc运算的批量算子数学上由erfc(x) 1 - erf(x)定义工程上以aclTensorList为输入输出形态通过两段式 aclnn 接口或图模式构图调用支持 FLOAT16/FLOAT32/BFLOAT16Kirin 平台仅前两者与 ND 格式输出要求连续内存输入支持非连续与空 Tensorshape 维度 0~8。其实现由算子 IR动态输入输出、Host 侧 Tiling累计偏移 多核切分与 Kernel 侧ForeachTriangle Erfc模板三层构成配合 CPU 仿真单测与 ATK 系统测试保证了计算的正确性。对于需要在 NPU 上批量计算高斯分布尾部概率、信号处理或概率统计场景的开发者可直接参考 示例代码 快速接入。赞分享人工智能算子库深度学习CANNAscend【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-nn点击查看免费下载相关推荐高性能中文语义嵌入实战指南基于bge-small-zh-v1.5的企业级应用架构设计高性能中文语义嵌入实战指南基于bge small zh v1.5的企业级应用架构设计 BGE Small ZH V1.5是BAAI北京智源人工智能研究院推人工智能算子库深度学习CANNAscend如何用 omc-doctor 检查 oh-my-claudecode 安装健康度并自动修复遗留 hooks 与版本漂移如何用 omc doctor 检查 oh my claudecode 安装健康度并自动修复遗留 hooks 与版本漂移 如果你在用 oh my claudec人工智能算子库深度学习CANNAscendCANN ops-math 数学算子 Erfc互补误差函数aclnnErfc 与 aclnnInplaceErfc 双段式接口使用指南CANN ops math 数学算子 Erfc互补误差函数aclnnErfc 与 aclnnInplaceErfc 双段式接口使用指南 导读 本文基于 C算子库人工智能CANN上一篇SmartOpenCV自定义绘制策略教程打造专属图像预览效果下一篇5 分钟跑通 lazygit终端里的 Git 可视化操作创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
