--- title: 基础矢量算子 description: "基于Ascend C方式实现基础矢量算子核函数的流程如下图所示。" url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/programug/Ascendcopdevg/atlas_ascendc_10_0033.html sourcePath: /source/zh/canncommercial/900/programug/Ascendcopdevg/atlas_ascendc_10_0033.html indexId: c8184b083719de04e715e3f4aaaed28919bdfc07c747f35c616804c32694b4db80 --- # 基础矢量算子 基于Ascend C方式实现基础矢量算子核函数的流程如下图所示。 图1 矢量算子核函数实现流程 - 算子分析：分析算子的数学表达式、输入、输出以及计算逻辑的实现，明确需要调用的Ascend C接口。 - 核函数定义：定义Ascend C算子入口函数。 - 根据 矢量编程范式 实现算子类：完成核函数的内部实现，包括3个基本任务：CopyIn，Compute，CopyOut。 下文以输入为half数据类型且shape的最后一维为32Bytes对齐、在单核上运行的、一次完成计算的Add算子为例，对上述步骤进行详细介绍。本样例中介绍的算子完整代码请参见基础Add算子样例(https://gitcode.com/cann/asc-devkit/tree/9.0.0/examples/01_simd_cpp_api/00_introduction/01_add/vector_add)。 #### 算子分析 算子分析具体步骤如下： 1. 明确算子的数学表达式及计算逻辑。 Add算子的数学表达式为： ``` z = x + y ``` 计算逻辑是：Ascend C提供的矢量计算接口(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0035.html)的操作元素都为LocalTensor(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0006.html)，输入数据需要先从外部存储（Global Memory）搬运进片上存储（Unified Buffer），然后使用计算接口完成两个输入参数相加，得到最终结果，再搬出到外部存储上。Ascend C Add算子的计算逻辑如下图所示。 图2 算子计算逻辑 2. 明确输入和输出。 - Add算子有两个输入：x与y；输出为z。 - 本样例中算子的输入支持的数据类型为half（float16），算子输出的数据类型与输入的数据类型相同。 - 算子输入支持的shape为（1，2048），输出shape与输入shape相同。 - 算子输入支持的 format 为：ND。 3. 确定核函数名称和参数。 - 您可以自定义核函数名称，本样例中核函数命名为vec_add_custom。 - 根据对算子输入输出的分析，确定核函数有3个参数x，y，z；x，y为输入在Global Memory上的内存地址，z为输出在Global Memory上的内存地址。 4. 确定算子实现所需接口。 - 实现涉及外部存储和内部存储间的数据搬运，查看Ascend CAPI参考中的数据搬运接口，需要使用 DataCopy(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0101.html) 来实现数据搬运。 - 本样例只涉及矢量计算的加法操作，查看Ascend CAPI参考中的 矢量计算(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0023.html) 接口，初步分析可使用基础算术Add接口 Add(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0035.html) 实现x+y。 - 使用 Queue队列(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0137.html) 管理计算中使用的Tensor数据结构，具体使用 EnQue(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0140.html) 、 DeQue(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0141.html) 等接口。 通过以上分析，得到Ascend C Add算子的设计规格如下： - 算子类型（OpType）：Add - 算子输入输出： **表1 Add算子输入输出规格** | name | shape | data type | format | | --- | --- | --- | --- | | x（输入） | (1, 2048) | half | ND | | y（输入） | (1, 2048) | half | ND | | z（输出） | (1, 2048) | half | ND | - 核函数名称：vec_add_custom - 使用的主要接口： - DataCopy：数据搬移接口 - Add：矢量基础算术接口 - EnQue、DeQue等接口：Queue队列管理接口 - 算子实现文件名称：vector_add.asc #### 核函数定义 根据核函数中介绍的规则进行核函数的定义。 1. 函数原型定义 本样例中，函数名为vector_add_custom（核函数名称可自定义），根据算子分析中对算子输入输出的分析，确定有3个参数x，y，z，其中x，y为输入内存，z为输出内存。根据核函数的规则介绍，函数原型定义如下所示：使用__global__函数类型限定符标识它是一个核函数，可以被<<<>>>调用；使用__vector__函数类型限定符标识该核函数在设备端aicore的Vector Core上执行；为方便起见，统一使用GM_ADDR宏修饰入参，GM_ADDR宏定义请参考核函数。 1 2 3 __global__ __vector__ void vector_add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { } 2. 调用算子类的Init和Process函数。 算子类的Init函数，完成内存初始化相关工作，Process函数完成算子实现的核心逻辑，具体介绍参见算子类实现。 1 2 3 4 5 6 7 __global__ __vector__ void vector_add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { AscendC::TPipe pipe; KernelAdd op; op.Init(x, y, z, &pipe); op.Process(); } 3. 根据核函数定义和调用章节，调用核函数时，除了需要传入参数x，y，z，还需要传入numBlocks（核函数执行的核数），nullptr（保留参数，设置为nullptr），stream（应用程序中维护异步操作执行顺序的stream）来规定核函数的执行配置。 1 vector_add_custom<<>>(xDevice, yDevice, zDevice); #### 算子类实现 根据上一节介绍，核函数中会调用算子类的Init和Process函数，本节具体讲解如何基于编程范式实现算子类。 根据矢量编程范式对Add算子的实现流程进行设计的思路如下，矢量编程范式请参考矢量编程范式，设计完成后得到的Add算子实现流程图参见图3 Add算子实现流程： - Add算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor xGm和yGm搬运至Local Memory，分别存储在xLocal，yLocal，Compute任务负责对xLocal，yLocal执行加法操作，计算结果存储在zLocal中，CopyOut任务负责将输出数据从zLocal搬运至Global Memory上的输出Tensor zGm中。 - CopyIn，Compute任务间通过VECIN队列inQueueX，inQueueY进行同步，Compute，CopyOut任务间通过VECOUT队列outQueueZ进行同步。 - 任务间交互使用到的内存、临时变量的内存统一使用 Pipe内存管理对象(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0108.html) 进行管理。 图3 Add算子实现流程 算子类中主要实现上述流程，包含对外开放的初始化Init函数和核心处理函数Process，Process函数中会对上图中的三个基本任务进行调用；同时包括一些算子实现中会用到的私有成员，比如上图中的GlobalTensor（xGm、yGm、zGm）和VECIN、VECOUT队列等。KernelAdd算子类具体成员如下： ``` class KernelAdd { public: __aicore__ inline KernelAdd() {} // 初始化函数，完成内存初始化相关操作 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, AscendC::TPipe* pipeIn){} // 核心处理函数，实现算子逻辑，调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作 __aicore__ inline void Process(){} private: // 搬入函数，完成CopyIn阶段的处理，被核心Process函数调用 __aicore__ inline void CopyIn(){} // 计算函数，完成Compute阶段的处理，被核心Process函数调用 __aicore__ inline void Compute(){} // 搬出函数，完成CopyOut阶段的处理，被核心Process函数调用 __aicore__ inline void CopyOut(){} private: AscendC::TPipe* pipe; // Pipe内存管理对象 AscendC::TQue inQueueX; // 输入数据Queue队列管理对象，TPosition为VECIN AscendC::TQue inQueueY; // 输入数据Queue队列管理对象，TPosition为VECIN AscendC::TQue outQueueZ; // 输出数据Queue队列管理对象，TPosition为VECOUT AscendC::GlobalTensor xGm; // 管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出 AscendC::GlobalTensor yGm; AscendC::GlobalTensor zGm; }; ``` 初始化函数主要完成以下内容： - 设置输入输出Global Tensor的Global Memory内存地址。 本样例中的分配方案是：数据整体长度TOTAL_LENGTH为1 * 2048，使用GlobalTensor类的SetGlobalBuffer(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0007.html)接口设定该核上Global Memory的起始地址以及长度。 1 xGm.SetGlobalBuffer((__gm__ half *)x, TOTAL_LENGTH); - 通过 Pipe内存管理对象(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0108.html) 为输入输出Queue分配内存。 比如，为输入x的Queue分配内存，可以通过如下代码段实现： 1 pipe->InitBuffer(inQueueX, 1, TOTAL_LENGTH * sizeof(half)) 具体的初始化函数代码如下： ``` constexpr int32_t TOTAL_LENGTH = 1 * 2048; // 数据总长 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, AscendC::TPipe* pipeIn) { pipe = pipeIn; // 设置Global Memory的起始地址以及长度 xGm.SetGlobalBuffer((__gm__ half*)x, TOTAL_LENGTH); yGm.SetGlobalBuffer((__gm__ half*)y, TOTAL_LENGTH); zGm.SetGlobalBuffer((__gm__ half*)z, TOTAL_LENGTH); // 通过Pipe内存管理对象为输入输出Queue分配内存 pipe->InitBuffer(inQueueX, 1, TOTAL_LENGTH * sizeof(half)); pipe->InitBuffer(inQueueY, 1, TOTAL_LENGTH * sizeof(half)); pipe->InitBuffer(outQueueZ, 1, TOTAL_LENGTH * sizeof(half)); } ``` 基于矢量编程范式，将核函数的实现分为3个基本任务：CopyIn，Compute，CopyOut。Process函数中通过如下方式调用这三个函数。 ``` __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } ``` 根据编程范式上面的算法分析，将整个计算拆分成三个Stage，用户单独编写每个Stage的代码，三阶段流程示意图参见图3，具体流程如下： 1. Stage1：CopyIn函数实现。 a. 使用 DataCopy(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0101.html) 接口将GlobalTensor数据拷贝到LocalTensor。 b. 使用 EnQue(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0140.html) 将LocalTensor放入VECIN的Queue中。 ``` __aicore__ inline void CopyIn() { // 从Que中为LocalTensor分配内存 AscendC::LocalTensor xLocal = inQueueX.AllocTensor(); AscendC::LocalTensor yLocal = inQueueY.AllocTensor(); // 将GlobalTensor数据拷贝到LocalTensor AscendC::DataCopy(xLocal, xGm, TOTAL_LENGTH); AscendC::DataCopy(yLocal, yGm, TOTAL_LENGTH); // LocalTensor放入VECIN的Queue中 inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal); } ``` 2. Stage2：Compute函数实现。 a. 使用 DeQue(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0141.html) 从VECIN中取出LocalTensor。 b. 使用Ascend C接口 Add(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0035.html) 完成矢量计算。 c. 使用 EnQue(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0140.html) 将计算结果LocalTensor放入到VECOUT的Queue中。 d. 使用 FreeTensor(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0139.html) 释放不再使用的LocalTensor。 ``` __aicore__ inline void Compute() { // 将Input从VECIN的Queue中取出 AscendC::LocalTensor xLocal = inQueueX.DeQue(); AscendC::LocalTensor yLocal = inQueueY.DeQue(); AscendC::LocalTensor zLocal = outQueueZ.AllocTensor(); // 调用Add算子进行计算 AscendC::Add(zLocal, xLocal, yLocal, TOTAL_LENGTH); // 将计算结果LocalTensor放入到VECOUT的Queue中 outQueueZ.EnQue(zLocal); // 释放LocalTensor inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); } ``` 3. Stage3：CopyOut函数实现。 a. 使用 DeQue(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0141.html) 接口从VECOUT的Queue中取出LocalTensor。 b. 使用 DataCopy(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0101.html) 接口将LocalTensor拷贝到GlobalTensor上。 c. 使用 FreeTensor(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0139.html) 将不再使用的LocalTensor进行回收。 ``` __aicore__ inline void CopyOut() { // 将计算结果从VECOUT的Queue中取出 AscendC::LocalTensor zLocal = outQueueZ.DeQue(); // 将计算结果从LocalTensor数据拷贝到GlobalTensor AscendC::DataCopy(zGm, zLocal, TOTAL_LENGTH); // 释放LocalTensor outQueueZ.FreeTensor(zLocal); } ```