一、概述本文档基于M300-P01开发板讲解Ascend C算子向量计算开发方法。首先系统阐释AI Core硬件架构核心知识要点，继而聚焦Transformer位置编码关键组件arange算子，通过完整工程实例深度演示向量化算子开发全流程。二、Ascend C简介 Ascend C是CANN针对算子开发场景推出的编程语言，原生支持C和C++标准规范，兼具开发效率和运行性能。基于Ascend C编写的算子程序，通过编译器编译和运行时调度，运行在昇腾AI处理器上。使用Ascend C，开发者可以基于昇腾AI硬件，高效的实现自定义的创新算法。使用Ascend C进行自定义算子开发的突出优势有： ●遵循C/C++编程规范，匹配用户开发习惯 ●自动并行调度，获得最优执行性能 ●结构化核函数编程，简化算子开发逻辑 ●CPU/NPU孪生调试，提升算子调试效率三、AI-CORE硬件架构介绍 1. 基本架构如下图所示，基于Ascend C开发的算子运行在AI Core上，对于初学者需要对硬件架构有基本认识，下文的编程模型基于硬件架构的抽象进行介绍，了解硬件架构能够帮助开发者更好的理解编程模型；对于需要完成高性能编程的深度开发者，更需要了解硬件架构相关知识，最佳实践中很多内容都以本章为基础进行介绍。 AI Core负责执行标量、向量和张量相关的计算密集型算子，包括三种基础计算单元：Cube（矩阵）计算单元、Vector（向量）计算单元和Scalar（标量）计算单元，同时还包含存储单元（包括硬件存储和用于数据搬运的搬运单元）和控制单元。 2. 计算单元计算单元是AI Core中提供强大算力的核心单元，包括三种基础计算单元：Cube（矩阵）计算单元、Vector（向量）计算单元和Scalar（标量）计算单元，完成AI Core中不同类型的数据计算。 Scalar Scalar负责各类型的标量数据运算和程序的流程控制。功能上可以看做一个小CPU，完成整个程序的循环控制、分支判断、Cube/Vector等指令的地址和参数计算以及基本的算术运算，并且可以通过在事件同步模块中插入同步符的方式来控制AI Core中其他执行单元的流水。相对于Host CPU，AI Core中的Scalar计算能力较弱，重点用于发射指令，所以在实际应用场景中应尽量减少Scalar计算，比如性能调优时尽量减少if/else等分支判断及变量运算。如下图所示：Scalar执行标量运算指令时，执行标准的ALU(Arithmetic Logic Unit)语句，ALU需要的代码段和数据段（栈空间）都来自于GM(global memory)，ICache(Instruction Cache)用于缓存代码段，缓存大小与硬件规格相关，比如为16K或32K，以2K为单位加载；DCache(Data Cache)用于缓存数据段，大小也与硬件规格相关，比如为16K，以cacheline（64Byte）为单位加载。考虑到核内访问效率最高，应尽量保证代码段和数据段被缓存在ICache和DCache，避免核外访问；同时根据数据加载单位不同，编程时可以考虑单次加载数据大小，来提升加载效率。例如在DCache加载数据时，当数据内存首地址与cacheline（64Byte）对齐时，加载效率最高。硬件提供L2Cache用于缓存访问GM的数据（包括代码段、数据段），以此加快访问速度，提高访问效率。核外L2Cache以cacheline为单位加载数据，根据硬件规格不同，cacheline大小不同（128/256/512Byte等）。 Vector Vector负责执行向量运算。向量计算单元执行向量指令，类似于传统的单指令多数据（Single Instruction Multiple Data，SIMD）指令，每个向量指令可以完成多个操作数的同一类型运算。如下图所示，向量计算单元可以快速完成两个FP16类型的向量相加或者相乘。向量指令支持多次迭代执行，也支持对带有间隔的向量直接进行运算。 Cube Cube计算单元负责执行矩阵运算，一次执行即可完成A矩阵（M K）与B矩阵（K N）的矩阵乘。如下图所示红色虚线框划出了Cube计算单元及其访问的存储单元，其中L0A存储左矩阵，L0B存储右矩阵，L0C存储矩阵乘的结果和中间结果。本文档演示的arange算子只涉及到Scalar和Vertor计算。 3. 存储单元 AI处理器中的计算资源要想发挥强劲算力，必要条件是保证输入数据能够及时准确地出现在计算单元中，需要精心设计存储系统，保证计算单元所需的数据供应。 AI Core中包含多级内部存储，AI Core需要把外部存储中的数据加载到内部存储中，才能完成相应的计算。AI Core的主要内部存储包括：L1 Buffer（L1缓冲区），L0 Buffer（L0缓冲区），Unified Buffer（统一缓冲区）等。为了配合AI Core中的数据传输和搬运，AI Core中还包含MTE（Memory Transfer Engine，存储转换引擎）搬运单元，在搬运过程中可执行随路数据格式/类型转换。 AI Core的内部存储列表如表1所示。存储单元描述 L1 Buffer L1缓冲区，通用内部存储，是AI Core内比较大的一块数据中转区，可暂存Cube计算单元需要反复使用的一些数据从而减少从总线读写的次数。 L0A Buffer / L0B Buffer Cube指令的输入。 L0C Buffer Cube指令的输出，但进行累加计算的时候，也是输入的一部分。 Unified Buffer 统一缓冲区，向量和标量计算的输入和输出。 BT Buffer BiasTable Buffer，存放Bias。 FP Buffer Fixpipe Buffer，存放量化参数、Relu参数等。搬运单元描述 MTE1 负责如下通路的数据搬运： ●L1->L0A/L0B ●L1->UB（只有耦合架构支持） ●L1->BT Buffer（只有分离架构支持） MTE2 负责如下通路的数据搬运： ●GM->{L1, L0A/B}，在该通路下，基于分形大小搬运，搬运时满足CacheLine（512B）对齐，性能更优。 ●GM->UB，基于CacheLine大小搬运性能更优。 MTE3 负责如下通路的数据搬运： ●UB -> GM ●UB -> L1（只有耦合架构支持） FixPipe 只有分离架构支持，负责如下通路的数据搬运，搬运过程中可以完成随路数据格式/类型转换： ●L0C->{GM/L1} ●L1->FP Buffer ●不同类型的AI处理器，存储单元大小不同，开发者可通过GetCoreMemSize接口获取。 ●所有通过搬运单元读写GM的数据都缺省被缓存在L2Cache，以此加快访问速度，提高访问效率。核外L2Cache以cacheline为单位加载数据，根据硬件规格不同，cacheline大小不同（128/256/512Byte等）。 4. 控制单元控制单元为整个计算过程提供了指令控制，负责整个AI Core的运行。控制单元/指令队列描述标量计算单元（Scalar） Scalar计算单元。矩阵运算队列（Cube Queue） Cube指令队列。同一个队列里的指令顺序执行，不同队列之间可以并行执行。向量运算队列（Vector Queue） Vector指令队列。同一个队列里的指令顺序执行，不同队列之间可以并行执行。存储转换队列（MTE Queue） MTE指令队列。同一个队列里的指令顺序执行，不同队列之间可以并行执行。事件同步模块（Event Sync）用于控制不同队列指令（也叫做不同指令流水）之间的依赖和同步的模块。 FP Buffer Fixpipe Buffer，存放量化参数、Relu参数等。多条指令从系统内存通过总线接口进入到ICache(Instruction Cache)中，后续的指令执行过程，根据指令的类型，有两种可能： ●如果指令是Scalar指令，指令会被Scalar单元直接执行。 ●其他指令会被Scalar单元调度到5个独立的分类队列（Vector指令队列、Cube指令队列、MTE1/MTE2/MTE3指令队列），然后再被对应执行单元执行。同一个指令队列中的指令是按照进入指令队列的顺序执行的，不同指令队列之间可以并行执行，通过多个指令队列的并行执行可以提升整体执行效率。对于并行执行过程中可能出现的数据依赖，通过事件同步模块插入同步指令来控制流水线的同步，提供PipeBarrier、SetFlag/WaitFlag两种指令，保证队列内部以及队列之间按照逻辑关系执行。 ●PipeBarrier本身是一条指令，用于在队列内部约束执行顺序（虽然指令是顺序执行，但并不意味着后一条指令开始执行时前一条指令执行结束）。PipeBarrier指令可以保证前序指令中所有数据读写全部完成，后序指令才开始执行。 ●SetFlag/WaitFlag为两条指令，在SetFlag/WaitFlag的指令中，可以指定一对指令队列的关系，表示两个队列之间完成一组“锁”机制，其作用方式为： ●SetFlag：当前序指令的所有读写操作都完成之后，当前指令开始执行，并将硬件中的对应标志位设置为1。 ●WaitFlag：当执行到该指令时，如果发现对应标志位为0，该队列的后续指令将一直被阻塞；如果发现对应标志位为1，则将对应标志位设置为0，同时后续指令开始执行。 Ascend C提供同步控制API，开发者可以使用这类API来自行完成同步控制。需要注意的是，通常情况下，开发者基于编程模型中介绍的编程模型和范式进行编程时不需要关注同步，编程模型帮助开发者完成了同步控制；使用编程模型和范式是我们推荐的编程方式，自行同步控制可能会带来一定的编程复杂度。四、工程化算子开发-arange Ascend C算子开发模式有两种，Kernel直调开发和工程化算子开发，为了方便部署，我们一般采用工程化算子开发模式。工程化算子开发是指基于自动生成的自定义算子工程完成算子实现、编译部署、单算子调用代码自动生成等一系列流程。本次样例arange算子，在Transformer的绝对位置编码中，arange负责生成连续位置索引序列。 1、算子分析 2、创建算子工程 CANN开发套件包中提供了自定义算子工程生成工具msOpGen，可基于算子原型定义输出算子工程：包括算子host侧代码实现文件、算子kernel侧实现文件以及工程编译配置文件等。使用msOpGen工具创建算子工程之前，需要参考环境准备章节安装驱动固件和CANN软件包，完成开发环境和运行环境的准备。使用msOpGen工具创建算子开发工程的步骤如下： 2.1 编写算子的原型定义json文件，用于生成算子开发工程。例如，ArangeCustom算子的json文件命名为arange_custom.json，文件内容如下 json [ { "op": "ArangeCustom", "language": "cpp", "input_desc": [ { "name": "start", "param_type": "required", "format": [ "ND" ], "type": [ "fp32" ] }, { "name": "end", "param_type": "required", "format": [ "ND" ], "type": [ "fp32" ] }, { "name": "step", "param_type": "required", "format": [ "ND" ], "type": [ "fp32" ] } ], "output_desc": [ { "name": "out", "param_type": "required", "format": [ "ND" ], "type": [ "fp32" ] } ] } ] 2.2 使用msOpGen工具生成算子的开发工程。下文仅针对关键参数进行解释，详细参数说明请参见算子工程创建（msOpGen）。 C++ ${INSTALL_DIR}/python/site-packages/bin/msopgen gen -i $HOME/sample/arange_custom.json -c ai_core-Ascend310P1 -lan cpp -out $HOME/sample/ArangeCustom ●${INSTALL_DIR}为CANN软件安装后文件存储路径，请根据实际环境进行替换。 ●-i：指定算子原型定义文件add_custom.json所在路径，请根据实际情况修改。 ●-c：ai_core- 代表算子在AI Core上执行，为昇腾AI处理器的型号。 ●-lan：参数cpp代表算子基于Ascend C编程框架，使用C/C++编程语言开发。 ●-out：生成文件所在路径，可配置为绝对路径或者相对路径，并且工具执行用户对路径具有可读写权限。若不配置，则默认生成在执行命令的当前路径。 2.3 命令执行完后，会在-out指定目录或者默认路径下生成算子工程目录，工程中包含算子实现的模板文件，编译脚本等，目录结构如下所示： C++ ├── ArangeCustom │ ├── build.sh │ ├── cmake │ │ ├── config.cmake │ │ ├── func.cmake │ │ ├── intf.cmake │ │ ├── makeself.cmake │ │ └── util │ ├── CMakeLists.txt │ ├── CMakePresets.json │ ├── framework │ │ ├── CMakeLists.txt │ │ └── tf_plugin │ ├── op_host │ │ ├── arange_custom.cpp │ │ ├── arange_custom_tiling.h │ │ └── CMakeLists.txt │ ├── op_kernel │ │ ├── arange_custom.cpp │ │ └── CMakeLists.txt │ └── scripts │ ├── help.info │ ├── install.sh │ └── upgrade.sh 工程目录中的op_kernel和op_host包含了算子的核心实现文件。op_kernel下存放kernel侧算子实现。op_host下存放host侧代码实现，包括算子原型定义、host侧tiling实现。其中kernel侧算子实现和host侧tiling实现在算子实现章节已经介绍了其核心的实现方法，在该章节会侧重于介绍接入CANN框架后的编程模式和API的使用。工程目录中的CMakePresets.json，用于开发者完成工程编译相关配置，之后即可进行编译部署。 3、算子实现 3.1 算子原型定义算子原型主要描述了算子的输入输出、属性等信息以及算子在AI处理器上相关实现信息，并关联tiling实现等函数。算子原型通过自定义的算子类来承载，该算子类继承自OpDef类。完成算子的原型定义等操作后，需要调用OP_ADD接口，传入算子类型（自定义算子类的类名），进行算子原型注册。下面是一个简单的Add算子原型定义和注册的例子。参考代码:op_host/arange_custom.cpp c++ namespace ops { class ArangeCustom : public OpDef { public: explicit ArangeCustom(const char *name) : OpDef(name) { this->Input("start") .ParamType(REQUIRED) .DataType({ge::DT_FLOAT}) .Format({ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND}) .Scalar(); this->Input("end") .ParamType(REQUIRED) .DataType({ge::DT_FLOAT}) .Format({ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND}) .Scalar(); this->Input("step") .ParamType(REQUIRED) .DataType({ge::DT_FLOAT}) .Format({ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND}) .Scalar(); this->Output("out") .ParamType(REQUIRED) .DataType({ge::DT_FLOAT}) .Format({ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND}); this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType); this->AICore().SetTiling(optiling::TilingFunc); this->AICore().AddConfig("ascend310p"); } }; 注册算子类型后，框架会根据算子类型获取算子注册信息，同时在编译和运行时按照一定的规则匹配算子实现文件名称和kernel侧核函数名称。为了保证正确匹配，算子类型、算子实现文件名称和核函数名称需要遵循如下定义规则。通常情况下，开发者只需要保证创建算子工程时原型定义json文件中算子类型op的参数值为大驼峰命名方式即可，工程创建后自动生成的代码即满足该规则。在手动编写算子原型定义和算子实现文件时需要按照如下规则定义。 ●算子类型需要采用大驼峰的命名方式，即采用大写字符区分不同的语义。 ●算子实现文件名称、核函数名称需相同，均为算子类型转换为下划线命名方式后的值。下文描述了通过算子类型转换成算子实现文件名称和核函数名称的过程： ●首字符的大写字符转换为小写字符。例如：Abc -> abc。 ●大写字符的前一个字符为小写字符或数字，则在大写字符前插一个下划线“_”，并将该字符转换为小写字符。例如：AbcDef -> abc_def。 ●大写字符前一个字符为大写字符且后一个字符是小写字符，则在大写字符前插一个下划线“_”，并将该字符转换为小写字符。例如：AbcAAc -> abc_a_ac。 ●其他大写字符转换为小写字符，小写字符保持不变。 3.2 Host侧tiling实现由于NPU中AI Core内部存储无法完全容纳算子输入输出的所有数据，需要每次搬运一部分输入数据进行计算然后搬出，再搬运下一部分输入数据进行计算，这个过程就称之为Tiling。切分数据的算法称为Tiling算法或者Tiling策略。根据算子的shape等信息来确定数据切分算法相关参数（比如每次搬运的块大小，以及总共循环多少次）的计算程序，称之为Tiling实现，也叫Tiling函数（Tiling Function）。由于Tiling实现中完成的均为标量计算，AI Core并不擅长，所以我们将其独立出来放在Host侧CPU上执行。为了最大化利用硬件资源，结合arange的输出长度进行了多级tiling切分，如下图： ●核间切分是为了利用多个AI core的并行处理 ●核间切分是考虑核内缓存资源限制，减少GM到UB的拷贝次数参考代码:op_host/arange_custom.cpp c++ namespace optiling { const uint32_t BLOCK_SIZE = 32; const uint32_t FLOAT_SIZE = 4; const uint32_t BLOCK_DIM = 2; #define DIVIDE_AND_ALIGN(size, split, align) \ ((((size) / (split)) + ((align) - 1)) & ~((align) - 1)) static ge::graphStatus TilingFunc(gert::TilingContext *context) { ArangeCustomTilingData tiling; uint32_t totalLength = context->GetOutputShape(0)->GetOriginShape().GetShapeSize(); ge::DataType dtype_out = context->GetOutputDesc(0)->GetDataType(); uint64_t ub_size; auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo()); ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ub_size); /*单次api计算大小：将ub 10等份后并按BLOCK_SIZE对齐*/ uint64_t ub_unit_size = DIVIDE_AND_ALIGN(ub_size, 10, BLOCK_SIZE); uint32_t singleCoreNum = (totalLength + BLOCK_DIM - 1) / BLOCK_DIM; uint32_t unitNum = ub_unit_size / FLOAT_SIZE; uint32_t unitLoops = singleCoreNum / unitNum; uint32_t tailNum = singleCoreNum - unitNum * unitLoops; if (tailNum > 0) unitLoops += 1; tiling.set_dtypeSize(FLOAT_SIZE); tiling.set_singleCoreNum(singleCoreNum); tiling.set_unitNum(unitNum); tiling.set_unitLoops(unitLoops); tiling.set_tailNum(tailNum); context->SetBlockDim(BLOCK_DIM); tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); size_t *currentWorkspace = context->GetWorkspaceSizes(1); currentWorkspace[0] = 0; return ge::GRAPH_SUCCESS; } } // namespace optiling 3.3 Kernel侧算子实现 Kernel实现即算子核函数实现，在Kernel函数内部通过解析Host侧传入的Tiling结构体获取Tiling信息，根据Tiling信息控制数据搬入搬出Local Memory的流程；通过调用计算、数据搬运、内存管理、任务同步API，实现算子逻辑。其核心逻辑基本上都为计算密集型任务，需要在NPU上执行。 arange算子kernel侧的设计思想是尽量使用vector计算，减少使用scalr计算：参考代码：op_kernel/arange_custom.cpp c++ #include "kernel_operator.h" constexpr int32_t BUFFER_NUM = 2; #define ALIGN_UP_32B_ELEMENTS(count, T) \ (((((count) sizeof(T)) + 31) / 32) (32 / sizeof(T))) template class KernelArange { public: __aicore__ inline KernelArange() {} __aicore__ inline void Init(GM_ADDR start, GM_ADDR end, GM_ADDR step, GM_ADDR out, uint32_t singleCoreNum, uint32_t unitNum, uint32_t unitLoops, uint32_t tailNum) { ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!"); startGm.SetGlobalBuffer((__gm__ float *)start); stepGm.SetGlobalBuffer((__gm__ float *)step); outGm.SetGlobalBuffer( (__gm__ float )out + singleCoreNum AscendC::GetBlockIdx(), singleCoreNum); pipe.InitBuffer(outQueue, BUFFER_NUM, unitNum * sizeof(float)); pipe.InitBuffer(temp1, unitNum * sizeof(float)); pipe.InitBuffer(temp2, unitNum * sizeof(float)); pipe.InitBuffer(temp3, unitNum * sizeof(float)); this->singleCoreNum = singleCoreNum; this->unitNum = unitNum; this->tailNum = tailNum; this->unitLoops = unitLoops; } __aicore__ inline void work_init() { TYPE_STEP step = stepGm.GetValue(0); TYPE_START start = startGm.GetValue(0) + singleCoreNum step AscendC::GetBlockIdx(); this->calc_init = temp1.Get (); this->calc_step = temp2.Get (); this->calc_temp = temp3.Get (); AscendC::CreateVecIndex(calc_temp, (TYPE_OUT)0, this->unitNum); AscendC::Duplicate(calc_init, start, this->unitNum); AscendC::Duplicate(calc_step, step, this->unitNum); AscendC::Mul(calc_step, calc_step, calc_temp, this->unitNum); AscendC::Add(calc_init, calc_init, calc_step, this->unitNum); AscendC::Duplicate(calc_temp, float(0.0), this->unitNum); this->offset_step_base = this->unitNum * step; AscendC::Duplicate(calc_step, this->offset_step_base, this->unitNum); } __aicore__ inline void Process() { /*初始化第一个UNIT序列值*/ work_init(); for (int32_t i = 0; i < this->unitLoops; i++) { if (i == this->unitLoops - 1 && this->tailNum > 0) { Compute(i, this->tailNum); CopyOut(i, this->tailNum); } else { Compute(i, this->unitNum); CopyOut(i, this->unitNum); } } } private: __aicore__ inline void Compute(int32_t iter, int32_t num) { uint32_t calc_num = ALIGN_UP_32B_ELEMENTS(num, float); AscendC::LocalTensor outLocal = outQueue.AllocTensor (); AscendC::Add(outLocal, this->calc_init, this->calc_temp, calc_num); AscendC::Add( this->calc_temp, this->calc_temp, this->calc_step, calc_num); outQueue.EnQue (outLocal); } __aicore__ inline void CopyOut(int32_t iter, int32_t num) { uint32_t copy_num = ALIGN_UP_32B_ELEMENTS(num, float); AscendC::LocalTensor outLocal = outQueue.DeQue (); // AscendC::DumpTensor(outLocal,iter, num); AscendC::DataCopy(outGm[iter * this->unitNum], outLocal, copy_num); outQueue.FreeTensor(outLocal); } AscendC::TPipe pipe; AscendC::TQue outQueue; AscendC::TBuf temp1, temp2, temp3; AscendC::GlobalTensor startGm; AscendC::GlobalTensor stepGm; AscendC::GlobalTensor outGm; AscendC::LocalTensor calc_init, calc_step, calc_temp; uint32_t singleCoreNum; uint32_t unitNum; uint32_t unitLoops; uint32_t tailNum; /*UNIT之间元素值差间隔*/ float offset_step_base; }; extern "C" __global__ __aicore__ void arange_custom(GM_ADDR start, GM_ADDR end, GM_ADDR step, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tiling_data, tiling); KernelArange op; op.Init(start, end, step, out, tiling_data.singleCoreNum, tiling_data.unitNum, tiling_data.unitLoops, tiling_data.tailNum); op.Process(); } 4、编译部署 4.1 算子工程编译算子kernel侧和host侧实现开发完成后，需要对算子工程进行编译，生成自定义算子安装包*.run，详细的编译操作包括： ●编译Ascend C算子kernel侧代码实现文件*.cpp，分为源码发布和二进制发布两种方式。 ○源码发布：不对算子kernel侧实现进行编译，保留算子kernel源码文件*.cpp。该方式可以支持算子的在线编译、通过ATC模型转换的方式编译算子的场景。 ○二进制发布：对算子kernel侧实现进行编译，生成描述算子相关信息的json文件*.json和算子二进制文件*.o。算子调用时，如果需要直接调用算子二进制，则使用该编译方式，比如通过单算子API调用的方式完成单算子的调用，PyTorch框架中单算子调用的场景，动态网络中调用算子的场景。 ●编译Ascend C算子host侧代码实现文件*.cpp、*.h。 ○将原型定义和shape推导实现编译成算子原型定义动态库libcust_opsproto_*.so，并生成算子原型对外接口op_proto.h。 ○将算子信息库定义编译成信息库定义文件*.json。 ○将Tiling实现编译成Tiling动态库liboptiling.so等。 ○基于算子原型定义，自动生成单算子API调用代码和头文件aclnn_*.h，并编译生成单算子API调用的动态库libcust_opapi.so。在算子工程目录下执行如下命令，进行算子工程编译 C++ bash build.sh 编译成功后，会在当前目录下创建build_out目录，并在build_out目录下生成自定义算子安装包custom_opp_ _ .run 4.2 算子包部署在自定义算子包所在路径下，执行如下命令，安装自定义算子包。 C++ ./build_out/custom_opp_openEuler_aarch64.run --install-path=$HOME/ops-bin 5、单算子调用单算子API调用方式，是指直接调用单算子API接口，基于C语言的API执行算子。算子工程创建完成后，基于工程代码框架完成算子原型定义、kernel侧算子实现、host侧tiling实现，通过工程编译脚本完成算子的编译部署，之后再进行单算子API的调用。 5.1 基本原理完成自定义算子编译后，会自动生成单算子API，可以直接在应用程序中调用。单算子API的形式一般定义为“两段式接口”，形如： C++ aclnnStatus aclnnArangeCustomGetWorkspaceSize(const aclScalar start, const aclScalar end, const aclScalar step, const aclTensor out, uint64_t workspaceSize, aclOpExecutor *executor); aclnnStatus aclnnArangeCustom(void workspace, uint64_t workspaceSize, aclOpExecutor executor, aclrtStream stream); 5.2 验证代码 ●验证代码目录: C++ ── AclNNInvocationNaive │ ├── build │ │ ├── CMakeCache.txt │ │ ├── CMakeFiles │ │ ├── cmake_install.cmake │ │ ├── execute_arange_op │ │ └── Makefile │ ├── CMakeLists.txt │ ├── gen_data.py │ ├── input │ │ └── input.bin │ ├── main.cpp │ ├── output │ │ ├── golden.bin │ │ └── output_y.bin │ ├── README.md │ ├── run.sh │ └── verify_result.py ●验证代码编译 ○ 设置cann的环境变量 bash export DDK_PATH=/usr/local/Ascend/ascend-toolkit/latest export NPU_HOST_LIB=/usr/local/Ascend/ascend-toolkit/latest/aarch64-linux/devlib ○设置自定义算子的环境变量 C++ source $HOME/ops-bin/vendors/customize/bin/set_env.bash ○样例执行 C++ bash run.sh ○预期输出 C++ test pass ##################################### INFO: you have passed the Precision!