基于Kernel直调工程的算子开发

环境准备

• CANN软件安装

  开发算子前，需要先准备好开发环境和运行环境，开发环境和运行环境的介绍和具体的安装步骤可参见《CANN 软件安装指南》。

• 环境变量配置

  安装CANN软件后，使用CANN运行用户进行编译、运行时，需要以CANN运行用户登录环境，执行source ${install_path}/set_env.sh命令设置环境变量。其中${install_path}为CANN软件的安装目录，例如：/usr/local/Ascend/ascend-toolkit。

算子分析

主要分析算子的数学表达式、输入输出的数量、Shape范围以及计算逻辑的实现，明确需要调用的Ascend C接口。下文以Add算子为例，介绍具体的分析过程。

1. 明确算子的数学表达式及计算逻辑。

   Add算子的数学表达式为：

   

   计算逻辑是：从外部存储Global Memory搬运数据至内部存储Local Memory，然后使用Ascend C计算接口完成两个输入参数相加，得到最终结果，再搬运到Global Memory上。

2. 明确输入和输出。
   • Add算子有两个输入：x与y，输出为z。
   • 本样例中算子输入支持的数据类型为half（float16），算子输出的数据类型与输入数据类型相同。
   • 算子输入支持的shape为（8，2048），输出shape与输入shape相同。
   • 算子输入支持的format为：ND。
3. 确定核函数名称和参数。
   • 本样例中核函数命名为add_custom。
   • 根据对算子输入输出的分析，确定核函数有3个参数x，y，z；x，y为输入参数，z为输出参数。
4. 确定算子实现所需接口。
   • 实现涉及外部存储和内部存储间的数据搬运，查看Ascend C API参考中的数据搬移接口，需要使用DataCopy来实现数据搬移。
   • 本样例只涉及矢量计算的加法操作，查看Ascend C API参考中的矢量计算接口矢量计算，初步分析可使用Add接口Add实现x+y。
   • 计算中使用到的Tensor数据结构，使用AllocTensor、FreeTensor进行申请和释放。
   • 并行流水任务之间使用Queue队列完成同步，会使用到EnQue、DeQue等接口。

通过以上分析，得到Ascend C Add算子的设计规格如下：

核函数开发

完成环境准备和初步的算子分析后，即可开始Ascend C核函数的开发。开发之前请先从矢量算子样例获取样例代码，以下样例代码在add_custom.cpp中实现。

本样例中使用多核并行计算，即把数据进行分片，分配到多个核上进行处理。Ascend C核函数是在一个核上的处理函数，所以只处理部分数据。分配方案是：假设共启用8个核，数据整体长度TOTAL_LENGTH为8 * 2048个元素，平均分配到8个核上运行，每个核上处理的数据大小BLOCK_LENGTH为2048个元素。下文的核函数，只关注长度为BLOCK_LENGTH的数据应该如何处理。

1. 首先，您需要根据核函数定义和调用中介绍的规则进行核函数的定义，并在核函数中调用算子类的Init和Process函数，算子类实现在后续步骤中介绍。

   1 2 3 4 5 6

   extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { KernelAdd op; op.Init(x, y, z); op.Process(); }

   • 使用__global__函数类型限定符来标识它是一个核函数，可以被<<<>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行。指针入参变量需要增加变量类型限定符__gm__，表明该指针变量指向Global Memory上某处内存地址。为了统一表达，使用GM_ADDR宏来修饰入参，GM_ADDR宏定义如下：

     1	#define GM_ADDR __gm__ uint8_t*

   • 算子类的Init函数，完成内存初始化相关工作，Process函数完成算子实现的核心逻辑。
2. 然后根据矢量编程范式实现算子类，本样例中定义KernelAdd算子类，其具体成员如下：

   1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24

   class KernelAdd { public: __aicore__ inline KernelAdd(){} // 初始化函数，完成内存初始化相关操作 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){} // 核心处理函数，实现算子逻辑，调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作 __aicore__ inline void Process(){} private: // 搬入函数，从Global Memory搬运数据至Local Memory，被核心Process函数调用 __aicore__ inline void CopyIn(int32_t progress){} // 计算函数，完成两个输入参数相加，得到最终结果，被核心Process函数调用 __aicore__ inline void Compute(int32_t progress){} // 搬出函数，将最终结果从Local Memory搬运到Global Memory上，被核心Process函数调用 __aicore__ inline void CopyOut(int32_t progress){} private: AscendC::TPipe pipe; //TPipe内存管理对象 AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY; //输入数据Queue队列管理对象，TPosition为VECIN AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ; //输出数据Queue队列管理对象，TPosition为VECOUT AscendC::GlobalTensor<half> xGm; //管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出 AscendC::GlobalTensor<half> yGm; AscendC::GlobalTensor<half> zGm; };

   内部函数的调用关系示意图如下：

   图2 核函数调用关系图
   

   由此可见除了Init函数完成初始化外，Process中完成了对流水任务“搬入、计算、搬出”的调用，开发者可以重点关注三个流水任务的实现。

3. 初始化函数Init主要完成以下内容：设置输入输出Global Tensor的Global Memory内存地址，通过TPipe内存管理对象为输入输出Queue分配内存。

   上文我们介绍到，本样例将数据切分成8块，平均分配到8个核上运行，每个核上处理的数据大小BLOCK_LENGTH为2048个元素。那么我们是如何实现这种切分的呢？

   每个核上处理的数据地址需要在起始地址上增加GetBlockIdx() * BLOCK_LENGTH（每个block处理的数据长度）的偏移来获取。这样也就实现了多核并行计算的数据切分。

   以输入x为例，x + BLOCK_LENGTH * GetBlockIdx()即为单核处理程序中x在Global Memory上的内存偏移地址，获取偏移地址后，使用GlobalTensor类的SetGlobalBuffer接口设定该核上Global Memory的起始地址以及长度。具体示意图如下。

   图3 多核并行处理示意图
   

   上面已经实现了多核数据的切分，那么单核上的处理数据如何进行切分？

   对于单核上的处理数据，可以进行数据切块（Tiling），在本示例中，仅作为参考，将数据切分成8块（并不意味着8块就是性能最优）。切分后的每个数据块再次切分成2块，即可开启double buffer，实现流水线之间的并行。

   这样单核上的数据（2048个数）被切分成16块，每块TILE_LENGTH（128）个数据。TPipe为inQueueX分配了两块大小为TILE_LENGTH * sizeof(half)个字节的内存块，每个内存块能容纳TILE_LENGTH（128）个half类型数据。数据切分示意图如下。

   图4 单核数据切分示意图
   

   具体的初始化函数代码如下：

   1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20

   #include "kernel_operator.h" constexpr int32_t TOTAL_LENGTH = 8 * 2048; // total length of data constexpr int32_t USE_CORE_NUM = 8; // num of core used constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // length computed of each core constexpr int32_t TILE_NUM = 8; // split data into 8 tiles for each core constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // separate to 2 parts, due to double buffer __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) { // get start index for current core, core parallel xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); // pipe alloc memory to queue, the unit is Bytes pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half)); }

4. 基于矢量编程范式，将核函数的实现分为3个基本任务：CopyIn，Compute，CopyOut。Process函数中通过如下方式调用这三个函数。

   1 2 3 4 5 6 7 8 9 10 11

   __aicore__ inline void Process() { // loop count need to be doubled, due to double buffer constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM; // tiling strategy, pipeline parallel for (int32_t i = 0; i < loopCount; i++) { CopyIn(i); Compute(i); CopyOut(i); } }

   1. CopyIn函数实现。
      1. 使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。
      2. 使用EnQue将LocalTensor放入VecIn的Queue中。

      1 2 3 4 5 6 7 8 9 10 11 12

      __aicore__ inline void CopyIn( int32_t progress) { // alloc tensor from queue memory AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>(); AscendC::LocalTensor<half> yLocal = inQueueY.AllocTensor<half>(); // copy progress_th tile from global tensor to local tensor AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH); AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH); // enque input tensors to VECIN queue inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal); }

   2. Compute函数实现。
      1. 使用DeQue从VecIn中取出LocalTensor。
      2. 使用Ascend C接口Add完成矢量计算。
      3. 使用EnQue将计算结果LocalTensor放入到VecOut的Queue中。
      4. 使用FreeTensor将释放不再使用的LocalTensor。

      1 2 3 4 5 6 7 8 9 10 11 12 13 14

      __aicore__ inline void Compute(int32_t progress) { // deque input tensors from VECIN queue AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>(); AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>(); AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>(); // call Add instr for computation AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH); // enque the output tensor to VECOUT queue outQueueZ.EnQue<half>(zLocal); // free input tensors for reuse inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); }

   3. CopyOut函数实现。
      1. 使用DeQue接口从VecOut的Queue中取出LocalTensor。
      2. 使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。
      3. 使用FreeTensor将不再使用的LocalTensor进行回收。

      1 2 3 4 5 6 7 8 9

      __aicore__ inline void CopyOut(int32_t progress) { // deque output tensor from VECOUT queue AscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>(); // copy progress_th tile from local tensor to global tensor AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH); // free output tensor for reuse outQueueZ.FreeTensor(zLocal); }

核函数运行验证

异构计算架构中，NPU（kernel侧）与CPU（host侧）是协同工作的，完成了kernel侧核函数开发后，即可编写host侧的核函数调用程序，实现从host侧的APP程序调用算子，执行计算过程。

除了上文核函数实现文件add_custom.cpp外，核函数的调用与验证还需要准备以下文件：

• 调用算子的应用程序：main.cpp。
• 输入数据和真值数据生成脚本文件：gen_data.py。
• 验证输出数据和真值数据是否一致的验证脚本：verify_result.py。
• 适用于CPU或NPU模式运行的算子编译工程文件：CMakeLists.txt。
• 编译运行算子的脚本：run.sh。

本文仅介绍main.cpp文件编写，其他内容您可以在矢量算子样例中直接获取。

1. host侧应用程序框架的编写。
   内置宏ASCENDC_CPU_DEBUG是区分运行CPU模式或NPU模式逻辑的标志，在同一个main函数中通过对ASCENDC_CPU_DEBUG宏定义的判断来区分CPU模式和NPU模式的运行程序。

   1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24

   #include "data_utils.h" #ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" extern void add_custom_do(uint32_t coreDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z); #else #include "tikicpulib.h" extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z); #endif int32_t main(int32_t argc, char* argv[]) { size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half uint32_t blockDim = 8; #ifdef ASCENDC_CPU_DEBUG // 用于CPU模式调试的调用程序 #else // NPU模式运行算子的调用程序 #endif return 0; }

2. 编写用于CPU调试的调用程序。
   完成算子核函数CPU模式运行验证的步骤如下：

   图5 CPU模式运行验证步骤
   

   GmAlloc、ICPU_RUN_KF、GmFree等接口说明请参考调测接口。

   1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16

   // 使用GmAlloc分配共享内存，并进行数据初始化 uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize); uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize); uint8_t* z = (uint8_t*)AscendC::GmAlloc(outputByteSize); ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize); ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize); // 调用ICPU_RUN_KF调测宏，完成核函数CPU模式调用 AscendC::SetKernelMode(KernelMode::AIV_MODE); ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug // 输出数据写出 WriteFile("./output/output_z.bin", z, outputByteSize); // 调用GmFree释放申请的资源 AscendC::GmFree((void *)x); AscendC::GmFree((void *)y); AscendC::GmFree((void *)z);

3. 编写NPU模式运行算子的调用程序。
   完成算子核函数NPU模式运行验证的步骤如下：

   图6 NPU模式运行验证步骤
   

   如下示例中的acl API使用方法请参考“acl API（C&C++）”章节。

   1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39

   // 初始化 CHECK_ACL(aclInit(nullptr)); // 运行管理资源申请 int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); aclrtStream stream = nullptr; CHECK_ACL(aclrtCreateStream(&stream)); // 分配Host内存 uint8_t *xHost, *yHost, *zHost; uint8_t *xDevice, *yDevice, *zDevice; CHECK_ACL(aclrtMallocHost((void**)(&xHost), inputByteSize)); CHECK_ACL(aclrtMallocHost((void**)(&yHost), inputByteSize)); CHECK_ACL(aclrtMallocHost((void**)(&zHost), outputByteSize)); // 分配Device内存 CHECK_ACL(aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); CHECK_ACL(aclrtMalloc((void**)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); CHECK_ACL(aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); // Host内存初始化 ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize); ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize); CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); // 用内核调用符<<<>>>调用核函数完成指定的运算,add_custom_do中封装了<<<>>>调用 add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice); CHECK_ACL(aclrtSynchronizeStream(stream)); // 将Device上的运算结果拷贝回Host CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST)); WriteFile("./output/output_z.bin", zHost, outputByteSize); // 释放申请的资源 CHECK_ACL(aclrtFree(xDevice)); CHECK_ACL(aclrtFree(yDevice)); CHECK_ACL(aclrtFree(zDevice)); CHECK_ACL(aclrtFreeHost(xHost)); CHECK_ACL(aclrtFreeHost(yHost)); CHECK_ACL(aclrtFreeHost(zHost)); // 去初始化 CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize());

4. 执行一键式编译运行脚本，编译和运行应用程序。
   脚本执行方式如下，<soc_version>表示算子运行的AI处理器型号，<run_mode>表示算子以cpu模式或npu模式运行。

   bash run.sh -r <run_mode> -v <soc_version>  

   1. 执行脚本前需要配置环境变量ASCEND_INSTALL_PATH，配置为CANN软件的安装路径，示例如下，请根据实际安装路径进行修改：

      export ASCEND_INSTALL_PATH=$HOME/Ascend/ascend-toolkit/latest

   2. 完成CPU模式和NPU模式运行验证。
      • CPU模式下执行如下命令，命令中的<soc_version>请替换为实际的AI处理器型号。

        bash run.sh -r cpu -v <soc_version>

      • NPU模式下执行如下命令，命令中的<soc_version>请替换为实际的AI处理器型号。

        bash run.sh -r npu -v <soc_version>

      当前使用numpy接口计算了输出数据和真值数据的绝对误差和相对误差，误差在容忍偏差范围内，视为精度符合要求，运行结果会输出"test pass"字样。

      

      AI处理器的型号<soc_version>请通过如下方式获取：

      • 非 Atlas A3 训练系列产品 / Atlas A3 推理系列产品 ：在安装昇腾AI处理器的服务器执行npu-smi info命令进行查询，获取Name信息。实际配置值为AscendName，例如Name取值为xxxyy，实际配置值为Ascendxxxyy。
      • Atlas A3 训练系列产品 / Atlas A3 推理系列产品 ：在安装昇腾AI处理器的服务器执行npu-smi info -t board -i id -c chip_id命令进行查询，获取Chip Name和NPU Name信息，实际配置值为Chip Name_NPU Name。例如Chip Name取值为Ascendxxx，NPU Name取值为1234，实际配置值为Ascendxxx_1234。

        其中：

        • id：设备id，通过npu-smi info -l命令查出的NPU ID即为设备id。
        • chip_id：芯片id，通过npu-smi info -m命令查出的Chip ID即为芯片id。

      该样例支持以下型号：

      • Atlas 推理系列产品
      • Atlas 训练系列产品
      • Atlas A2 训练系列产品 / Atlas 800I A2 推理产品 /A200I A2 Box 异构组件

接下来的引导

如果您对教程中的多核并行、流水编程等概念不了解，导致阅读过程有些吃力，可以参考编程模型学习基本概念，再来回顾本教程；如果您已经了解相关概念，并跑通了该样例，您可以参考矢量编程了解Ascend C矢量编程中的更多细节。