ATVC 模板库开发指南：用 Ascend C Templates for Vector Compute 快速构建 Vector 算子

ATVC 模板库开发指南用 Ascend C Templates for Vector Compute 快速构建 Vector 算子【免费下载链接】atvcATVCAscend C Templates for Vector Compute是为基于Ascend C开发的典型Vector算子封装的一系列模板头文件的集合可帮助用户快速开发典型Vector算子。项目地址: https://gitcode.com/cann/atvcATVCAscend C Templates for Vector Compute是一个基于 Ascend C API 构建的 C 模板头文件集合面向在昇腾 AI 处理器上开发典型 Vector 算子的场景将算子开发中的 Kernel 数据搬入搬出、资源申请与通用 Tiling 计算等底层逻辑封装为可复用的模板模块。阅读本文后你将掌握 ATVC 的环境搭建、编译安装与 UT/样例验证流程理解 Host 层与 Kernel 层的分层设计及OpTraits/Param/Policy三大公共数据结构并了解 Elementwise、Reduce、Broadcast、Pool 四类算子模板及其组合的开发方式。1. ATVC 是什么ATVC 是一个用 Ascend C API 搭建的 C 模板头文件集合旨在帮助用户快速开发 Ascend C 典型 Vector 算子。它将 Ascend C Vector 算子开发流程中的计算实现解耦成可自定义的模块内部封装实现了 Kernel 数据搬入搬出等底层通用操作及通用 Tiling 计算形成了高效的算子开发模式。相比传统 Ascend C 算子开发方式利用 ATVC 搭建的 Vector 算子可提升开发效率 3-5 倍该提升比例来自项目官方说明用户只需选择匹配的模板并完成核心计算逻辑即可完成算子 Kernel 侧开发ATVC 还内置了每个模板库对应的通用 Tiling 计算实现可省去手写 Tiling 的开发量在保证不错性能表现的同时极大提升算子开发效率。从源码结构看整个库以include/atvc.h作为对外总入口头文件include/atvc.h它统一收拢了common/atvc_opdef.hOpTraits 定义、elewise/host/elewise_host.h、reduce/host/reduce_host.h、broadcast/host/broadcast_host.h等 Host 侧头文件并在非 Host 侧编译条件下引入 Kernel 侧模板头文件elewise_op_template.h、reduce_op_template.h、broadcast_op_template.h等。用户可直接通过导入 include 路径下的对外头文件include/atvc.h使能模板库能力。2. 分层架构与核函数调用关系ATVC 将 Vector 算子开发流程中的可定制化模块抽象出了 Host 层和 Kernel 层基本概念如下Host 层在 CPU Host 侧执行提供 Tiling 计算与策略分派的 API它根据实际数据场景帮助用户计算出较优的数据搬运等运行态参数。Kernel 层利用 Ascend C API 搭建的一系列 Vector 算子模板类屏蔽了算子开发中用户无需感知的数据搬入搬出以及资源申请等固定模块并将核心计算模块开放给用户定义。基于如上分层结构ATVC 中一个核函数的实现与调用的关系如下图所示ATVC 框架提供的模板及接口用黄色表示支持开发自定义的模块用蓝色表示从调用链路看Host 侧通过内置/自定义 Tiling 计算 API 产出运行态参数经由可选的策略分派 API内置/自定义 Adapter选择最佳 Kernel 实现Kernel 侧通过内置 Kernel 算子模板 计算模板内置 Compute/自定义 Compute提供 Vector 算子核心逻辑并由算子原型模板OpTraits描述编译态算子原型。这一设计与 include/atvc.h 中 Host/Kernel 头文件的分段包含方式一致#ifndef __NPU_HOST__分支才引入 Kernel 模板保证了 Host 与 Kernel 侧编译的隔离。3. 代码目录结构ATVC 代码目录结构如下├── build.sh # 项目工程编译脚本 ├── cmake # 项目工程编译目录 ├── CMakeLists.txt # 编译配置文件 ├── docs # 项目文档介绍 ├── examples # ATVC 样例 ├── include # 项目公共头文件 ├── README.md ├── scripts # 项目脚本文件存放目录 ├── test # UT测试工程目录各目录职责与仓库内容对应如下include/按算子类别组织模板头文件包含elewise/、reduce/、broadcast/、pool/四个子目录每个子目录下再按common公共定义、hostTiling 与调度 API、kernel算子模板类与计算模板划分另有common/存放atvc_opdef.h、type_list.h、dtype_utils.h等公共基础头文件。examples/提供add、add_with_broadcast、addcmul、broadcast_to、reduce_sum、relu_with_reduce_sum、sinh_custom、tanh_grad、edge等样例以及ops_aclnnAclNN 单算子调用样例与ops_pytorchPyTorch 扩展样例等更贴近真实使用场景的工程。test/ut/UT 测试工程目录block/下包含test_add.cpp、test_reduce_sum.cpp等对应各算子样例的用例tiling/下包含test_atvc_tiling.cpp等 Tiling 用例可通过根目录build.sh脚本批量执行。scripts/存放打包、安装相关的脚本如scripts/package/package.py及各 shell 脚本。cmake/项目工程编译目录含模块查找Finddlog.cmake、Findsecurec.cmake、第三方依赖gtest 等与打包配置。4. 支持的模板与数据类型ATVC 当前提供四类算子模板均支持int32_t与float两种数据类型规格限制如下算子模板数据类型规格限制说明Elementwiseint32_t、floatReduceint32_t、float当前只支持 4 维以内的 reduce 计算Broadcastint32_t、float当前只支持 2 维以内对齐场景的计算Poolint32_t、float只支持 2 维以内场景的计算各类模板的 Host 层与 Kernel 层 API 名称如下算子模板Host 层 APIKernel 层模板类ElementwiseCalcEleWiseTilingEleWiseOpTemplateReduceCalcReduceTilingReduceOpTemplateBroadcastCalcBroadcastTilingBroadcastOpTemplatePoolCalcPoolTilingPoolOpTemplate4.1 算子类型说明Elementwise 类算子通常是指对张量进行元素级别操作的函数或方法包括但不限于加、减、乘、除及指数、对数、三角函数等数学函数。特点是逐元素进行计算操作不改变输入数据的形状。常见的有 Add、Sub、Exp、Log、Sin、Sqrt 等。Reduce 类算子通常是对张量中的元素进行归约操作的算子常用来求和、求平均值等可指定某几个维度进行归约计算也可将所有元素归约计算为一个标量。常见的有 ReduceSum求和、ReduceMean求平均值、ReduceProduct累乘、ReduceMax、ReduceMin、ReduceAny、ReduceAll 等。Broadcast 类算子用于在张量形状不一致时实现张量间的逐元素运算。例如张量 A 的 Shape 为 (1, 5)张量 B 的 Shape 为 (3, 5)为完成 C A B需依据广播规则将 A 由 (1, 5) 扩展至 (3, 5)通过在长度为 1 的维度上复制数据使两个张量的形状对齐。Pool 类算子常用于神经网络的池化层达到降采样、特征提取等功能。例如输入 Shape 为 (4, 4)池化窗口为 (2, 2)输入数据为[[1,2,3,4],[5,6,7,8],[9,10,11,12],[13,14,15,16]]步幅为 2、不使用填充时最大池化结果为[[6,8],[14,16]]。4.2 三大公共数据结构ATVC 核函数定义与调用涉及三个公共数据概念OpTraits编译态参数参考 C 模板元编程的type_list实现提供ATVC::OpInputs、ATVC::OpOutputs、ATVC::OpTemps三个模板结构体分别描述算子的计算输入、计算输出与临时资源支持 C 基础类型作为不定长模板参数。三者组合成ATVC::OpTraits。完整定义见 include/common/atvc_opdef.h。例如一个c a b的 Add 算子可描述为using AddInputs ATVC::OpInputsfloat, float; // 两个输入类型均为float using AddOutputs ATVC::OpOutputsfloat; // 一个输出类型为float using AddTemps ATVC::OpTemps; // 无需临时buffer模板参数为空 using AddOpTraits ATVC::OpTraitsAddInputs, AddOutputs, AddTemps; // Add算子的计算原型描述Param运行态参数框架提供ATVC::EleWiseParam、ATVC::ReduceParam、ATVC::BroadcastParam以及 Pool 场景的ATVC::PoolParam结构体描述算子内部调度的 Tiling 数据和其他资源变量。Param 作为 Host 侧 Tiling API 的输出传入 Kernel 层算子模板在运行时指导数据循环搬运与调度计算。Policy模板策略的编译态参数ATVC::ReducePolicy、ATVC::BroadcastPolicy是 Kernel 层对部分算子模板的拓展描述对应算子模板类在不同场景的实例化实现。它由 Tiling API 计算出并在策略分派 API如ATVC::Host::ReduceAdapter里将运行态的 Policy 结果转化为模板参数调用该场景下的最佳模板实现。5. 快速入门若希望快速体验项目可按照以下流程依次完成环境准备、源码下载、编译安装、UT 测试与样例运行验证。5.1 依赖项ATVC 支持源码编译进行源码编译前需确保系统满足以下要求AI 处理器型号Ascend 910C、Ascend 910B。安装依赖gawkpython 3.7.0gcc 7.3.0cmake 3.16.0gtest可选仅执行 UT 时依赖gtest 需下载源码后自行编译安装参考步骤如下mkdir temp cd temp # 在gtest源码根目录下创建临时目录并进入 cmake .. -DCMAKE_CXX_FLAGS-fPIC -D_GLIBCXX_USE_CXX11_ABI0 make make install # root用户安装gtest # sudo make install # 非root用户安装gtest5.2 环境准备安装驱动与固件运行态依赖模板算子运行时必须安装驱动与固件安装指导详见 CANN 软件安装指南见 快速入门文档 中的官方安装指引。安装 CANN toolkit 包根据实际产品型号和环境架构获取Ascend-cann-toolkit_${cann_version}_linux-${arch}.run安装包后执行# 需要确保安装目录权限至少为755 # 确保安装包具有可执行权限 chmod x Ascend-cann-toolkit_${cann_version}_linux-${arch}.run # 安装命令 ./Ascend-cann-toolkit_${cann_version}_linux-${arch}.run --full --force --install-path${install_path}参数说明${cann_version}CANN 包版本号。${arch}CPU 架构如 aarch64、x86_64。${install_path}指定安装路径缺省--install-path时使用默认路径安装。若使用 root 用户安装安装完成后相关软件存储在/usr/local/Ascend/路径下若使用非 root 用户安装存储在$HOME/Ascend/路径下。环境变量设置示例默认路径root 用户安装source /usr/local/Ascend/cann/set_env.sh默认路径非 root 用户安装source $HOME/Ascend/cann/set_env.sh指定路径安装source ${install_path}/cann/set_env.sh5.3 源码下载git clone https://gitcode.com/cann/atvc.git用户可直接通过导入源码路径下 include 路径中的对外头文件include/atvc.h使能模板库能力。在进行算子开发之前可以通过下述 UT 测试和样例运行验证确保当前环境准备已完备。5.4 编译安装环境准备好后在开源仓根目录执行以下命令编译生成可部署的安装包bash build.sh --pkg编译完成后在仓库根目录的build_out目录下生成安装包例如cann-atvc_${version}_linux-${arch}.run。执行生成的安装包完成安装# 确保安装包具有可执行权限 chmod x cann-atvc_${version}_linux-${arch}.run # 安装命令 ./cann-atvc_${version}_linux-${arch}.run --full --quiet --install-path${install_path}参数说明${version}表示安装包版本号如 8.5.0${arch}表示 CPU 架构如 aarch64、x86_64${install_path}指定安装路径缺省时安装至默认路径。5.5 UT 测试可选在开源仓根目录执行下列命令之一将依次批跑 test 目录下的用例得到结果日志用于看护编译是否正常bash build.sh -u或bash build.sh --utest如需显示测试覆盖率依赖 lcov 1.14bash build.sh --utest --cov5.6 样例运行验证开发者调用 ATVC 实现自定义算子开发后可通过单算子调用的方式验证算子功能。本仓提供部分算子实现及其调用样例具体参考 examples 目录下的样例。以add算子样例为例# 从根目录进入到样例的目录 cd examples # 编译执行命令 bash run_examples.sh add若提示如下信息则说明算子运行成功精度比较通过。更详细的用例执行流程参阅 样例执行说明... Accuracy verification passed. Sample add passed!从 examples/run_examples.sh 的实现看样例脚本依赖bisheng编译器通过-I ${ATVC_HOME_DIR}/include引入 ATVC 模板头文件并以--npu-archdav-2201指定昇腾 NPU 架构成功执行后输出Sample ${TEST_NAME} passed!。以 examples/sinh_custom/sinh_custom.cpp 为例样例内部先调用ATVC::Host::CalcEleWiseTilingSinhOpTraits(eleNum, param)计算运行态参数再以SinhCustomSinhOpTraitsparam.tilingData.blockNum, nullptr, stream(xDevice, yDevice, param)直调核函数最后通过VerifyResults完成 golden 数据比对并打印Accuracy verification passed.。6. 编程指南要点速览完整开发教程见 编程指南。以下提炼四类模板开发的关键路径与核心约束便于快速定位。6.1 Elementwise 算子开发自定义 Elementwise 算子按以下顺序组装模块定义计算模板 → 将计算模板类传入 Kernel 层算子模板完成核函数实现 → 定义 Kernel 层算子入口 API内部实例化计算模板类。以 Sinh 算子y (exp(x) - exp(-x)) / 2为例完整样例见 examples/sinh_custom/sinh_custom.cpp#include atvc.h // 包含所有模板及API的总入口头文件 // 描述算子的输入输出以及临时计算资源 using SinhOpTraits ATVC::OpTraitsATVC::OpInputsfloat, ATVC::OpOutputsfloat, ATVC::OpTempsfloat, float; // 传入编译态参数ATVC::OpTraits templateclass Traits struct SinhComputeFunc { templatetypename T, typename U __aicore__ inline void operator()(AscendC::LocalTensorT x, AscendC::LocalTensorT y, AscendC::LocalTensorU tempBuffer1, AscendC::LocalTensorU tempBuffer2) { uint32_t tiledCnt y.GetSize(); // 进行单次基块计算的元素个数 AscendC::Muls(tempBuffer1, x, static_castT(-1), tiledCnt); // tempBuffer1 -1 * x AscendC::Exp(tempBuffer1, tempBuffer1, tiledCnt); // tempbuffer1 exp(-x) AscendC::Exp(tempBuffer2, x, tiledCnt); // tempbuffer2 exp(x) AscendC::Sub(y, tempBuffer2, tempBuffer1, tiledCnt); // y exp(x) - exp(-x) AscendC::Muls(y, y, static_castT(0.5), tiledCnt); // y (e^(x) - e^(-x)) / 2 } }; templateclass OpTraits __global__ __aicore__ void SinhCustom(GM_ADDR x, GM_ADDR y, ATVC::EleWiseParam param) { KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); // 控制算子执行时只启动Vector核 auto op ATVC::Kernel::EleWiseOpTemplateSinhComputeFuncOpTraits(); op.Run(x, y, param); // OpTraits内部的ATVC::OpTemps将由EleWiseOpTemplate内部申请资源开发无需关注 }计算模板类的定义约束实例化时固定传入ATVC::OpTraits类型结构体作为模板参数。必须重载公有仿函数__aicore__ inline void operator()由ATVC::Kernel::EleWiseOpTemplate在计算阶段调用。仿函数形式参数需按照ATVC::OpInputs、ATVC::OpOutputs、ATVC::OpTemps声明的顺序填入其他标量参数放在最后按需传入。核函数定义约束必须预留一个GM_ADDR类型的形参用于传入ATVC::EleWiseParam运行态参数。核函数内部必须加入KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);标注算子执行时只启动 Vector 核。必须初始化ATVC::Kernel::EleWiseOpTemplate变量并调用其Run(Args... args)接口实现数据调度运算。若OpTraits固定如算子输入类型不变核函数定义与调用可简化为extern C形式直接传入固定的SinhOpTraits。Host 侧通过ATVC::Host::CalcEleWiseTilingOpTraits(totalCnt, param)计算 Tiling其内部依据EleWiseTilingData含tailBlockCnt、tailElemCnt、numPerBlock、tiledCnt、blockNum与EleWiseParam含totalCnt、nBufferNum指导数据搬运核心切分逻辑可参考 include/elewise/host/elewise_host.h。6.2 Reduce 算子开发Reduce 模板算子内部根据数据大小、Shape 与 Reduce 轴完成不同计算调度的实现各场景抽象为ATVC::ReducePolicy。由于涉及多核数据结果同步与核内分块对齐计算ATVC 提供内置计算模板如ATVC::ReduceSumCompute完整代码见 include/reduce/kernel/reduce_sum_compute.h。Reduce 计算模板必须实现以下公有 APICompute(...)计算单数据基块的 Reduce 结果。UpdateCache(...)计算单 UB 内不同数据基块的计算结果。ReduceBetweenUB(...)计算多核之间及同一核内的多次 UB 结果。GetPaddingValue()返回非对齐场景下不参与计算的尾部数据填充值。完整样例见 examples/reduce_sum/reduce_sum.cpp其关键调用链为CalcReduceTilingReduceOpTraits(shape, dim, policy, param)计算 Tiling 与 Policy →ReduceOpAdapterReduceOpTraits(...)完成 workspace 申请、Policy 转编译态参数并实例化ReduceCustomOpTraits, ATVC::REDUCE_POLICYx核函数。核函数需加入KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);标注多核控制指令场景。从 include/reduce/kernel/reduce_op_template.h 的模板定义可见ReduceOpTemplate除ReduceCompute与SelectReducePolicy外还预留了PreCompute、PostCompute模板参数用于组合算子扩展。6.3 Broadcast 算子开发Broadcast 模板算子内部根据数据类型、输入/输出 Shape 完成轴上数据扩充各计算调度场景抽象为ATVC::BroadcastPolicy。完整样例见 examples/broadcast_to/broadcast_to.cpp其Compute接口需完成两种场景的计算AB 场景输入src为 Shape(dimA, 1)扩充到 Shape(dimA, dimB) 的dst。BA 场景输入src为 Shape(1, dimA)扩充到 Shape(dimB, dimA) 的dst。内置计算模板见 include/broadcast/kernel/broadcast_compute.hHost 侧CalcBroadcastTiling的实现参考 include/broadcast/host/broadcast_host.h。6.4 组合算子开发ATVC 支持基于模板参数扩展的组合算子Broadcast 组合支持 Elementwise Broadcast、Broadcast Elementwise、Elementwise Broadcast Elementwise完整样例见 examples/add_with_broadcast/add_with_broadcast.cpp 与 examples/add_with_broadcast/post_compute_add_of_broadcast.hReduce 组合支持 Elementwise ReduceSum、ReduceSum Elementwise、Elementwise ReduceSum Elementwise完整样例见 examples/relu_with_reduce_sum/relu_with_reduce_sum.cpp 与 examples/relu_with_reduce_sum/post_compute_relu_with_reduce_sum.h。组合算子中用户定义的OpTraits为组合算子的整体定义核函数内部再据此拆分出 Elementwise 与 Broadcast/Reduce 各自的OpTraits。以 AddWithBroadcast 为例对应源码 examples/add_with_broadcast/add_with_broadcast.cppBroadcast 的输入是组合算子的第一个输入Broadcast 的输出是组合算子的临时资源后置 Elementwise 计算模板PostComputeAddOfBroadcast除重载operator()外还需定义SetArgs接收向量参数与SetParam接收标量参数两个接口。BroadcastOpTemplate通过HAS_PRE_COMPUTE/HAS_POST_COMPUTE编译期判断决定用PreCompute替换 CopyIn、用PostCompute替换 CopyOut并支持SetParam完成融合场景如融合激活系数的晚期参数注入。6.5 Pool 算子开发Pool 算子的模块交互与 Elementwise 模板类似但计算模板类必须额外设置TILE_LAYOUT与TILE_PADDING两个编译期常量描述基本块大小与上下左右外扩元素个数。完整样例见 examples/edge/edge.cpp// 计算基本块宽高 宽需要32B对齐 static constexpr ATVC::Layout2Dim TILE_LAYOUT{16, 16}; // tile块上下左右padding的设置 left/right需要32B对齐 static constexpr ATVC::PoolTilePadding TILE_PADDING{8, 8, 1, 1};其中Layout2Dim含width、heightPoolTilePadding含left、right、up、down。Host 侧CalcPoolTiling依据总 layout 与 tile layout 计算切分块数、启动核数及每核处理块数详细算法见 include/pool/host/pool_host.h模板类定义见 include/pool/kernel/pool_op_template.h。7. 调试调优功能7.1 OpTraits 校验接口用户可通过DebugCheck()接口在 Host 侧校验不同模板的 OpTraits无需额外开关限制namespace ATVC { namespace Host { template typename OpTraits, ATVC::TemplateType templateType bool DebugCheck(); } }TemplateType枚举含ELE_WISE、REDUCE、BROADCAST三类校验项如下模板类型OpTraits 校验项ELE_WISE输入输出非空REDUCE输入输出个数均为 1输入输出数据类型相同BROADCAST输入输出个数均为 1输入输出数据类型相同使用示例using AddOpTraits ATVC::OpTraitsATVC::OpInputsfloat, float, ATVC::OpOutputsfloat; ATVC::Host::DebugCheckAddOpTraits, ATVC::TemplateType::ELE_WISE();完整的DebugCheck调用样例可参考 tanh_grad 算子 examples/tanh_grad/tanh_grad.cpp。7.2 调试调优运行模式样例执行脚本 examples/run_examples.sh 支持可选入参--run-mode选择调试调优模式--run-modedebug_printDFX 信息打印模式打开 kernel 侧模板内置关键节点信息打印和异常退出打印内部映射-DATVC_DEBUG_MODE1。--run-modeprofilingProfiling 性能采集模式打开性能数据采集内部映射-DATVC_DEBUG_MODE2并以msprof --ai-coreon ...启动采集。未设置--run-mode默认模式正常上板无 kernel 侧 DFX 打印、未开启 profiling。DFX 信息打印格式为[日志级别(ERROR/INFO)]:[ATVC][Module] (可选:[CopyIn/CopyOut等])其中ATVC标识模板库内置打印Module标识如EleWise、Reduce、Broadcast、Common等模块。模板内部提供的打印接口为ATVC::Kernel::DebugPrintf普通算子开发用户无需关注仅在需要修改或扩展模板功能时使用。为增加 profiling 采集稳定性建议在开启 profiling 时重复多次调用 kernel 以消除抖动。7.3 Tiling 超参调优ATVC 的 Tiling API 均支持传入超参结构体进行性能探索ElementwiseEleWiseTilingHyperParam可调singleCoreBaseLine单核数据量基线范围 [256, 128*1024]默认 512、ubSizeLimitThresholdUB 内存使用上限比例[0.5, 0.96]默认 0.95、nBufferNum并行流水 buffer 数[1, 2]默认 2、splitDataShape单核内数据量 3 个分段节点默认 {1024, 321024, 641024}、dataSplitFactor4 个数据段切分系数均需在 [1, 32]默认 {4, 4, 8, 6}、rsvLiveCnt预留 UB 空间比例[0, 1]默认 0。通用切分算法为按totalCnt / singleCoreBaseLine计算blockNum最小 1、最大为平台 vectorCore 上限按 UB 上限与单元素内存计算ubLimitCnttiledCnt依据avgElePerBlock所处数据段以切分系数确定且需不超过ubLimitCnt、不小于 32 并做 32 元素对齐再据此计算tailBlockCnt、numPerBlock、tailElemCnt。ReduceReduceTilingHyperParam可调basicBlock基本块内存大小不超过 UB 内存 1/3192K 内存建议 48K-54K默认 48*1024、maxInnerAAR 切轴内 A 轴最大数据量[128, 256]默认 128、balanceThreshHold多核均衡阈值水平[0.8, 0.95]默认 0.85。超参修改调用示例Elementwiseusing AddOpTraits ATVC::OpTraitsATVC::OpInputsfloat, float, ATVC::OpOutputsfloat; int32_t eleNum 8 * 1024; ATVC::EleWiseParam param; ATVC::Host::EleWiseTilingHyperParam hyperParam; hyperParam.singleCoreBaseLine 1024; if (!ATVC::Host::CalcEleWiseTilingAddOpTraits(eleNum, param, hyperParamhyperParam)) { printf([ERROR]: Calculate eleWise tiling failed.\n); return -1; };未传入自定义超参时使用各超参结构体的默认值。8. 文档与相关信息快速入门快速体验项目的简易教程环境准备、编译执行、本地验证等。编程指南使用 ATVC 实现算子开发的完整教程公共数据结构、四类模板开发、组合算子、调试调优。贡献指南安全声明许可证ATVC 项目由 CANN 社区开源维护以模板化、模块化的方式大幅降低典型 Vector 算子的开发门槛适合在 Ascend 910B/910C 平台上进行算子原型验证与高性能开发。上手时建议按“环境准备 → 样例验证 → 单模板开发 → 组合算子开发”的路径循序渐进并结合 examples 下的完整样例与 test/ut 中的测试用例对照学习。【免费下载链接】atvcATVCAscend C Templates for Vector Compute是为基于Ascend C开发的典型Vector算子封装的一系列模板头文件的集合可帮助用户快速开发典型Vector算子。项目地址: https://gitcode.com/cann/atvc创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考