• 开发者分享 | Ascend C算子开发及单算子调用

    本文分享自《AscendC算子开发及单算子调用》，作者：goldpancake。

    笔者在阅读Ascend C官方文档的过程中发现，对于初学者来说，尤其是第一次接触异构编程思想的初学者，有部分内容是无需特别关注的，例如算子工程的相关的CmakeLists.txt，以及单算子调用的一些通用工具类文件。同时，在环境配置的过程中，也发现了一些需要注意的地方，特此记录备忘。

    1 环境准备

    笔者的硬件及系统环境如下：

    • 操作系统：openEuler release 20.03 (LTS-SP3)

    • 设备：Ascend 910

    开发环境需要准备三个run包，分别是驱动、固件和cann-toolkit开发套件，笔者这里使用当前的最新版本CANN开发套件，版本号为7.0.RC1.alpha003，并在昇腾社区下载好对应驱动和固件的run包。

    1.1 安装流程

    上述准备的三个包，按照驱动 -> 固件 -> CANN开发套件包的顺序来安装。

    首先安装驱动，执行如下命令：

    /path/to/Ascend-hdk-910-npu-driver_23.0.rc2_linux-aarch64.run --full --install-for-all

    注意：笔者使用root用户进行安装，以full模式执行run包，并加上install-for-all选项来为所有用户安装。

    接下来安装固件，执行如下命令：

    /path/to/Ascend-hdk-910-npu-firmware_6.4.12.1.241.run --full

    驱动和固件都安装完成后，最好重启一次系统：

    reboot

    重启完成后，安装CANN开发套件包：

    path/to/Ascend-cann-toolkit_7.0.RC1.alpha003_linux-aarch64.run --full --install-for-all

    安装完成后，开发环境就准备好了。

    1.2 安装过程中可能的问题

    笔者在安装过程中，遇到了一个问题，很蠢，但值得注意。

    问题的表现是，在按照上述的流程安装好开发环境之后，除root用户外的其他普通用户使用msopgen工具生成算子工程时，出现了权限不足的问题。但因为加上了install-for-all选项，所以不应该是CANN包的权限问题。然后又查看msopgen的代码发现，该工具将python解释器指定为了root用户下的conda环境中的解释器。

    1. #!/root/miniconda3/bin/python3
    2. # coding=utf-8
    3. """
    4. Function:
    5. This file mainly involves main function of op generation module.
    6. Copyright Information:
    7. Huawei Technologies Co., Ltd. All Rights Reserved © 2020
    8. """

    原来是root用户下的conda配置为了默认激活base环境，笔者安装时没有注意这一点，导致在CANN包安装的过程中，选择到了conda环境下的python解释器，这样一来，其他用户肯定是没有权限的。在关闭base环境重新安装CANN包后，问题解决。

    2 算子开发流程

    至此，环境准备好后，开始正式的算子开发步骤。

    2.1 算子工程配置文件

    CANN包中提供了一个自动生成算子工程的工具msopgen，该工具可以通过一个json配置文件来生成完整的算子工程，具体的编写方式请参考Ascend C官方文档。

    这里以sinh算子为例，该算子是一元操作，所以只需要一个输入，且输出形状与输入形状一致。根据该特征来编写json文件，为了贴合Ascend C官方建议的编程范式，将文件命名为sinh_custom.json。为了简洁，这里我们只实现一种数据类型的操作。

    1. [
    2.     {
    3.         "op": "SinhCustom",
    4.         "language": "cpp",
    5.         "input_desc": [
    6.             {
    7.                 "name": "x",
    8.                 "param_type": "required",
    9.                 "format": [
    10.                     "ND"
    11.                 ],
    12.                 "type": [
    13.                     "fp16"
    14.                 ]
    15.             }
    16.         ],
    17.         "output_desc": [
    18.             {
    19.                 "name": "y",
    20.                 "param_type": "required",
    21.                 "format": [
    22.                     "ND"
    23.                 ],
    24.                 "type": [
    25.                     "fp16"
    26.                 ]
    27.             }
    28.         ]
    29.     }
    30. ]
     

    2.2 生成算子工程
     

    创建一个文件夹用作算子工程目录，使用msopgen工具执行如下命令来生成算子工程。
     

    1. mkdir /path/to/SinhCustom
    2. /path/to/msopgen gen -i /path/to/sinh_custom.json -c ai_core-Ascend910 -lan cpp -out /path/to/SinhCustom
     

    命令行会输出类似如下的信息：
     

    1. 2023-10-07 14:58:42 (942445) - [INFO] Start to generate AI Core operator files.
    2. 2023-10-07 14:58:42 (942445) - [INFO] Start to parse the ir template:/path/to/SinhCustom/sinh_custom.json
    3. 2023-10-07 14:58:42 (942445) - [INFO] Start to parse the op: SinhCustom
    4. 2023-10-07 14:58:42 (942445) - [INFO] Start to parse the input_desc: x
    5. 2023-10-07 14:58:42 (942445) - [INFO] Start to parse the output_desc: y
    6. 2023-10-07 14:58:42 (942445) - [WARNING] The "attr" value is invalid or no "attr" exists in the map.
    7. 2023-10-07 14:58:42 (942445) - [INFO] Start to check the type and format between the inputs/outputs in IR template.
    8. 2023-10-07 14:58:42 (942445) - [INFO] Start to generate a new project.
    9. 2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/cmake/config.cmake generated successfully.
    10. 2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/op_host/sinh_custom_tiling.h generated successfully.
    11. 2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/op_host/sinh_custom.cpp generated successfully.
    12. 2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/op_kernel/sinh_custom.cpp generated successfully.
    13. 2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/framework/tf_plugin/tensorflow_sinh_custom_plugin.cc generated successfully.
    14. 2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/framework/tf_plugin/CMakeLists.txt generated successfully.
    15. 2023-10-07 14:58:42 (942445) - [INFO] Generation completed.
     

    此时会发现指定的输出目录只已经生成了一系列的算子工程文件。
     

    1. SinhCustom
    2. ├── build.sh
    3. ├── cmake
    4. ├── CMakeLists.txt
    5. ├── CMakePresets.json # 这个配置项需要修改
    6. ├── framework
    7. ├── op_host
    8. │   ├── CMakeLists.txt
    9. │   ├── sinh_custom.cpp # 算子host侧核心逻辑
    10. │   └── sinh_custom_tiling.h # 算子tiling结构体定义
    11. ├── op_kernel
    12. │   ├── CMakeLists.txt
    13. │   └── sinh_custom.cpp # 算子kernel侧核心逻辑
    14. ├── scripts
    15. └── sinh_custom.json # 笔者此处将工程配置文件和算子工程目录放在了一起
     

    我们只需要专注于上述带有注释的几个文件即可。
     

    此处先修改与算子核心逻辑无关的配置项CMakePresets.json，官方文档中也描述的非常清楚，只需要将ASCEND_CANN_PACKAGE_PATH配置项修改为实际的CANN包安装路径即可。在root用户下安装的默认路径为/usr/local/Ascend/ascend-toolkit/latest。
     

    以上将所有无关算子逻辑的内容修改完毕，接下来就可以专注于算子开发了。
     

     

    2.3 算子逻辑开发
     

    官方文档中推荐先实现kernel侧的逻辑，但笔者有一些不同的看法。我推荐先实现算子tiling结构体的定义与具体策略，这样做的好处是，可以提前将tiling策略所需的变量确定下来，并且借助于CANN包只提供的一系列宏，这一过程并不需要很大的工作量。在实现kernel侧逻辑的过程中，这些变量将有助于思考数据在逻辑核上如何具体分配和执行，当然这只是笔者的观点，可以根据自己的编程习惯作调整。
     

     

    2.3.1 tiling结构体定义及策略实现
     

    首先确定tiling过程中所需的变量，参考官方样例，需要定义整块、尾块的个数及其中的元素个数，还需要定义最小对齐单位。op_host/sinh_custom_tiling.h代码如下：
     

    1. #ifndef SINH_CUSTOM_TILING_H // 头文件保护记得加上，自动生成的文件中不包含
    2. #define SINH_CUSTOM_TILING_H
    3. #include "register/tilingdata_base.h"
    4.  
    5. namespace optiling
    6. {
    7.   BEGIN_TILING_DATA_DEF(TilingData)
    8.   TILING_DATA_FIELD_DEF(uint32_t, formerNum);    // 整块个数
    9.   TILING_DATA_FIELD_DEF(uint32_t, tailNum);      // 尾块个数
    10.   TILING_DATA_FIELD_DEF(uint32_t, formerLength); // 整块内元素个数
    11.   TILING_DATA_FIELD_DEF(uint32_t, tailLength);   // 尾块内元素个数
    12.   TILING_DATA_FIELD_DEF(uint32_t, alignNum);     // 最小对齐单位，元素个数
    13.   END_TILING_DATA_DEF;
    14.  
    15.   REGISTER_TILING_DATA_CLASS(SinhCustom, TilingData)
    16. }
    17.  
    18. #endif
     

     然后在op_host/sinh_custom.cpp中实现具体的tiling策略，代码如下： 

     

    1. namespace optiling
    2. {
    3.     constexpr uint32_t BLOCK_DIM = 24;                        // 划分核心数量
    4.     constexpr uint32_t SIZE_OF_HALF = 2;                      // 数据类型的字节数
    5.     constexpr uint32_t BLOCK_SIZE = 32;                       // 昇腾设备上的数据block为32字节
    6.     constexpr uint32_t ALIGN_NUM = BLOCK_SIZE / SIZE_OF_HALF; // 最小对齐单位
    7.     static ge::graphStatus TilingFunc(gert::TilingContext *context)
    8.     {
    9.  
    10.         TilingData tiling;
    11.         uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize();
    12.         context->SetBlockDim(BLOCK_DIM);
    13.  
    14.         // 使输入向上对齐
    15.         uint32_t totalLengthAligned = ((totalLength + ALIGN_NUM - 1) / ALIGN_NUM) * ALIGN_NUM;
    16.         // 计算整块和尾块个数
    17.         uint32_t formerNum = (totalLengthAligned / ALIGN_NUM) % BLOCK_DIM;
    18.         uint32_t tailNum = BLOCK_DIM - formerNum;
    19.         // 计算整块和尾块的元素个数
    20.         uint32_t formerLength = ((totalLengthAligned / BLOCK_DIM + ALIGN_NUM - 1) / ALIGN_NUM) * ALIGN_NUM;
    21.         uint32_t tailLength = (totalLengthAligned / BLOCK_DIM / ALIGN_NUM) * ALIGN_NUM;
    22.  
    23.         // 设置tiling参数
    24.         tiling.set_formerNum(formerNum);
    25.         tiling.set_tailNum(tailNum);
    26.         tiling.set_formerLength(formerLength);
    27.         tiling.set_tailLength(tailLength);
    28.         tiling.set_alignNum(ALIGN_NUM);
    29.  
    30.         // 以下为固定写法，不用纠结
    31.         tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
    32.         context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
    33.         context->SetTilingKey(1);
    34.         size_t *currentWorkspace = context->GetWorkspaceSizes(1);
    35.         currentWorkspace[0] = 0;
    36.  
    37.         return ge::GRAPH_SUCCESS;
    38.     }
    39. }
     

    2.3.2 kernel侧实现
     

    有了上述实现的tiling策略，我们就可以根据数据划分的逻辑来确定kernel侧的具体实现。根据官方推荐的矢量编程范式，我们可以先将算子类的框架写出来，再慢慢填充内容。在op_kernel/sinh_custom.cpp中写出算子类框架。
     

    1. using namespace AscendC; // 记得开启AscendC命名空间
    2. constexpr int32_t BUFFER_NUM = 2; // TQue的缓冲数量，此处开启双Buffer
    3.  
    4. class KernelSinh
    5. {
    6. public:
    7.     __aicore__ inline KernelSinh() {} // 类构造函数，无须任何代码
    8.     __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,   // 初始化函数的参数为输入、输出
    9.                                 uint32_t formerNum, uint32_t tailNum, // 以及上面定义的一系列tiling参数
    10.                                 uint32_t formerLength, uint32_t tailLength,
    11.                                 uint32_t alignNum) { /* TODO */ }
    12.     __aicore__ inline void Process() { /* TODO */ }
    13.  
    14. private:
    15.     __aicore__ inline void CopyIn() { /* TODO */ }
    16.     __aicore__ inline void Compute() { /* TODO */ }
    17.     __aicore__ inline void CopyOut() { /* TODO */ }
    18.  
    19. private:
    20.     /* TODO */
    21. };
     

    第一步，分析算子类的私有数据成员。
     

    首先一定需要的是用来管理内存的Tpipe，同时需要输入输出分别对应的TQue和GlobalTensor，同时每个逻辑核还需要直到当前处理的数据个数，所以需要一个变量tileLength来确定分片大小。
     

    第二步，分析算子。
     

    公式：
    y=sinh(x)=ex−e−x2.0" role="presentation" style="text-align: center; position: relative;">y=sinh(x)=ex−e−x2.0
     

    可以观察到，我们需要计算两个中间结果，分别是$e^{\bf x}$和$e^{-{\bf x}}$，所以需要相应的数据结构来存放这两个中间结果，Ascend C提供的TBuf可以很好的承担这一责任。
     

    至此我们就将算子类需要的私有数据成员确定了下来。
     

    1. TPipe pipe;                                      // 用于操作队列
    2. TBuf tempBuf;              // 存放中间结果
    3. TQue inQueueX;   // 输入队列
    4. TQue outQueueY; // 输出队列
    5. GlobalTensor xGm;                       // 输入数据对应的GM内存空间
    6. GlobalTensor yGm;                       // 输出数据对应的GM内存空间
    7. uint32_t tileLength;                             // 每个逻辑核需要知道分片数据个数
     

    第三步，完善算子类的初始化函数Init()。
     

    在该函数中我们需要为GlobalTensor分配内存，并初始化相应的TQue，同时需要针对某些变量做合法性判断。
     

    1. __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,
    2.                             uint32_t formerNum, uint32_t tailNum,
    3.                             uint32_t formerLength, uint32_t tailLength,
    4.                             uint32_t alignNum)
    5. {
    6.     if (GetBlockIdx() < formerNum)
    7.     {
    8.         // 处理整块逻辑
    9.         this->tileLength = formerLength;
    10.         xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + formerLength * GetBlockIdx(), formerLength);
    11.         yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + formerLength * GetBlockIdx(), formerLength);
    12.     }
    13.     else
    14.     {
    15.         // 处理尾块逻辑
    16.         this->tileLength = tailLength;
    17.         xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + formerLength * formerNum + tailLength * (GetBlockIdx() - formerNum), tailLength);
    18.         yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + formerLength * formerNum + tailLength * (GetBlockIdx() - formerNum), tailLength);
    19.     }
    20.  
    21.     ASSERT(alignNum != 0 && "align num can not be zero!");
    22.     pipe.InitBuffer(inQueueX, BUFFER_NUM, (((this->tileLength + alignNum - 1) / alignNum) * alignNum) * sizeof(half));
    23.     pipe.InitBuffer(outQueueY, BUFFER_NUM, (((this->tileLength + alignNum - 1) / alignNum) * alignNum) * sizeof(half));
    24. }
     

    第四步，完成算子最核心的部分：根据矢量编程范式实现算子计算逻辑。
     

    1. __aicore__ inline void CopyIn()
    2. {
    3.     LocalTensor xLocal = inQueueX.AllocTensor();
    4.     DataCopy(xLocal, xGm, this->tileLength); // GM -> LM
    5.     inQueueX.EnQue(xLocal);
    6. }
    7. __aicore__ inline void Compute()
    8. {
    9.     LocalTensor xLocal = inQueueX.DeQue();
    10.     LocalTensor yLocal = outQueueY.AllocTensor();
    11.     pipe.InitBuffer(tempBuf, this->tileLength * sizeof(DTYPE_X));
    12.     LocalTensor tempLocal = tempBuf.Get(this->tileLength);
    13.     // 计算exp(x)
    14.     Exp(yLocal, xLocal, this->tileLength);
    15.     // 计算-x
    16.     half nagOne(-1.0);
    17.     Muls(tempLocal, xLocal, nagOne, this->tileLength);
    18.     // 计算exp(-x)
    19.     Exp(tempLocal, tempLocal, this->tileLength);
    20.     // 计算exp(x)-exp(-x)
    21.     Sub(yLocal, yLocal, tempLocal, this->tileLength);
    22.     // 计算最终结果
    23.     half denominator(0.5);
    24.     Muls(yLocal, yLocal, denominator, this->tileLength);
    25.     outQueueY.EnQue(yLocal);
    26.     inQueueX.FreeTensor(xLocal);
    27. }
    28. __aicore__ inline void CopyOut()
    29. {
    30.     LocalTensor yLocal = outQueueY.DeQue();
    31.     DataCopy(yGm, yLocal, this->tileLength); // LM -> GM
    32.     outQueueY.FreeTensor(yLocal);
    33. }
     

    实现的具体细节与接口可以参考Ascend C官方文档。
     

    第五步，将Process()函数补全，并完善核函数。
     

    1. __aicore__ inline void Process()
    2. {
    3.     CopyIn();
    4.     Compute();
    5.     CopyOut();
    6. }
    7.  
    8. extern "C" __global__ __aicore__ void
    9. sinh_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
    10. {
    11.     GET_TILING_DATA(tiling_data, tiling);
    12.     KernelSinh op;
    13.     op.Init(x, y,
    14.             tiling_data.formerNum, tiling_data.tailNum,
    15.             tiling_data.formerLength, tiling_data.tailLength,
    16.             tiling_data.alignNum);
    17.     if (TILING_KEY_IS(1))
    18.     {
    19.         op.Process();
    20.     }
    21. }
     

    至此就完成了kernel侧的实现。
     

     

    2.3.3 host侧实现
     

    我们回到op_host/sinh_custom.cpp，关于类型推导函数，这个算子输入输出的形状一致。msopgen生成的算子工程中，默认即为输入输出形状一致，所以无须改动。如果在写其他复杂算子的时候，需要仔细分析数据形状的变化。关于算子原型注册，也无须改动。
     

    现在就完成了整个算子的逻辑，可以执行build.sh来验证有没有编译时错误，若没有错误则可以进行运行时验证。
     

     

    3 核函数调用
     

    笔者直接将官方的核函数调用样例拿来做了一些修改，需要修改的地方如下。
     

    1. kernel_invocation
    2. ├── cmake
    3. ├── CMakeLists.txt
    4. ├── data_utils.h
    5. ├── input
    6. ├── main.cpp # 需要修改
    7. ├── output
    8. ├── run.sh # 需要修改
    9. ├── add_custom.cpp # 替换为自己的算子实现
    10. ├── add_custom.py # 需要修改
    11. └── verify_result.py # 添加的代码，用于验证结果
     

    首先，将官方样例中的add_custom.cpp替换为自己实现的kernel侧算子，笔者这里的名称为sinh_custom.cpp。同时为了CPU侧调试，需要添加一个核函数的包装函数，代码如下。
     

    1. #ifndef __CCE_KT_TEST__
    2. void sinh_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y)
    3. {
    4.     sinh_custom<<>>(x, y);
    5. }
    6. #endif
     

    注意：为了快速验证逻辑，在核函数验证过程中未使用动态tiling，所以没有之前提到的那些tiling参数。
     

    然后是sinh_custom.py，官方样例中是add_custom.py，这里修改文件名称，因为后面的run.sh中是通过算子文件名来调用这一python脚本的。
     

    由于本算子只需要一个输入向量，所以只生成一个input数据，然后修改golden数据的生成方式，调用numpy中与算子功能相同的函数来计算，注意数据类型，代码如下。
     

    1. import numpy as np
    2.  
    3. def gen_golden_data_simple():
    4.     np.random.seed(42)
    5.     input_x = np.random.randn(8, 2048).astype(np.float16)
    6.     golden = np.sinh(input_x).astype(np.float16)
    7.     print(f'-----------------------{input_x[0][0]}')
    8.     input_x.tofile("./input/input_x.bin")
    9.     golden.tofile("./output/golden.bin")
    10.  
    11. if __name__ == "__main__":
    12.     gen_golden_data_simple()
     

    main.cpp中要调整相应的内存申请等操作，只需要一个input，CPU侧调试和NPU侧调试的代码都需要修改，具体如下。
     

    1. #include 
    2.  
    3. #include "data_utils.h"
    4. #ifndef __CCE_KT_TEST__
    5. #include "acl/acl.h"
    6. extern void sinh_custom_do(uint32_t coreDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y);
    7. #else
    8. #include "tikicpulib.h"
    9. extern "C" __global__ __aicore__ void sinh_custom(GM_ADDR x, GM_ADDR y);
    10. #endif
    11.  
    12. int32_t main(int32_t argc, char *argv[])
    13. {
    14.   size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    15.   size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);
    16.   uint32_t blockDim = 8;
    17.  
    18. #ifdef __CCE_KT_TEST__
    19.   uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputByteSize);
    20.   uint8_t *y = (uint8_t *)AscendC::GmAlloc(outputByteSize);
    21.  
    22.   ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
    23.  
    24.   AscendC::SetKernelMode(KernelMode::AIV_MODE);
    25.   ICPU_RUN_KF(sinh_custom, blockDim, x, y);
    26.  
    27.   WriteFile("./output/output_y.bin", y, outputByteSize);
    28.  
    29.   AscendC::GmFree((void *)x);
    30.   AscendC::GmFree((void *)y);
    31. #else
    32.   CHECK_ACL(aclInit(nullptr));
    33.   aclrtContext context;
    34.   int32_t deviceId = 0;
    35.   CHECK_ACL(aclrtSetDevice(deviceId));
    36.   CHECK_ACL(aclrtCreateContext(&context, deviceId));
    37.   aclrtStream stream = nullptr;
    38.   CHECK_ACL(aclrtCreateStream(&stream));
    39.  
    40.   uint8_t *xHost, *yHost;
    41.   uint8_t *xDevice, *yDevice;
    42.   CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputByteSize));
    43.   CHECK_ACL(aclrtMallocHost((void **)(&yHost), outputByteSize));
    44.   CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    45.   CHECK_ACL(aclrtMalloc((void **)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    46.  
    47.   ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
    48.   CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
    49.  
    50.   sinh_custom_do(blockDim, nullptr, stream, xDevice, yDevice);
    51.   CHECK_ACL(aclrtSynchronizeStream(stream));
    52.  
    53.   CHECK_ACL(aclrtMemcpy(yHost, outputByteSize, yDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
    54.   WriteFile("./output/output_y.bin", yHost, outputByteSize);
    55.  
    56.   CHECK_ACL(aclrtFree(xDevice));
    57.   CHECK_ACL(aclrtFree(yDevice));
    58.   CHECK_ACL(aclrtFreeHost(xHost));
    59.   CHECK_ACL(aclrtFreeHost(yHost));
    60.  
    61.   CHECK_ACL(aclrtDestroyStream(stream));
    62.   CHECK_ACL(aclrtDestroyContext(context));
    63.   CHECK_ACL(aclrtResetDevice(deviceId));
    64.   CHECK_ACL(aclFinalize());
    65. #endif
    66.   return 0;
    67. }
     

     

    原样例中的验证方式是求md5和，但由于核函数中调用了Exp、Muls等API，所以精度可能会有损失，不适合用md5sum的方式来验证。这里就需要引入新的文件verify_result.py，这里使用了numpy.isclose函数来进行验证，这也是官方单算子API调用的结果验证方式。
     

    1. import sys
    2. import math
    3. import numpy as np
    4.  
    5. def data_compare(file1, file2,file3):
    6.     input1 = np.fromfile(file1, dtype=np.float16)
    7.     print("input1: ", input1)
    8.     golden = np.fromfile(file2, dtype=np.float16)
    9.     output = np.fromfile(file3, dtype=np.float16)
    10.     print("output: ", output)
    11.     print("-------------golden is :")
    12.     print("golden: ", golden)
    13.  
    14.     different_element_results = np.isclose(
    15.         output, golden,
    16.         rtol=5e-2,
    17.         atol=1e-3,
    18.         equal_nan=True)
    19.     different_element_indexes = np.where(
    20.         different_element_results != np.array((True,)))[0]
    21.     if different_element_indexes.size == 0:
    22.         print("result correct!")
    23.     else:
    24.         print("result error!")
    25.     return 0 if different_element_indexes.size == 0 else 1
    26.  
    27. if __name__ == '__main__':
    28.     intput_file1 = sys.argv[1]
    29.     golden_file = sys.argv[2]
    30.     output_file = sys.argv[3]
    31.     cmp_result = data_compare(intput_file1, golden_file, output_file)
    32.  
    33.     if (cmp_result == 0):
    34.         sys.exit(0)
    35.     else:
    36.         sys.exit(1)
     

    最后是修改run.sh脚本，需要修改的只有最后验证结果的部分。
     

    原样例的验证方式是md5sum：
     

    echo "md5sum: ";md5sum output/*.bin
     

    修改为调用脚本判断：
     

    echo "result verification: " python3 verify_result.py ./input/input_x.bin ./output/golden.bin ./output/output_y.bin
     

     

    4 单算子API调用
     

    单算子调用是通过自动生成的两段式API来执行的，为了快速验证，同样是将官方样例中的单算子API调用样例拿来做了一些修改。需要修改的几处关键代码如下。
     

    1. aclnn_online_model
    2. ├── build
    3. ├── inc
    4. ├── README.md
    5. ├── run
    6. │   └── out
    7. │       ├── execute_sinh_op
    8. │       ├── result_files
    9. │       └── test_data
    10. │           ├── config
    11. │           └── data
    12. │               ├── generate_data.py # 生成测试数据脚本，需要修改
    13. ├── run.sh # 需要修改
    14. ├── scripts
    15. │   └── verify_result.py # 调整验证方式，例如相对和绝对误差参数等
    16. └── src
    17.     ├── CMakeLists.txt # 需要修改
    18.     ├── common.cpp
    19.     ├── main.cpp # 需要修改
    20.     ├── operator_desc.cpp
    21.     └── op_runner.cpp # 需要修改
     

    具体细节如下。
     

    generate_data.py中，按照算子来修改测试数据生成方式。本算子需要half类型的测试数据，故代码改为：
     

    1. import numpy as np
    2. a = np.random.randn(8, 2048).astype(np.float16)
    3. a.tofile('input_0.bin')
     

    verify_result.py中，根据实际读取的输入和输出，利用np.isclose来进行比较，该函数详细用法参考numpy官方文档。
     

    1. import sys
    2. import math
    3. import numpy as np
    4.  
    5. def data_compare(file1, file2):
    6.     input1 = np.fromfile(file1, dtype=np.float16)
    7.     print("input1: ", input1)
    8.     golden = np.sinh(input1).astype(np.float16)
    9.     output = np.fromfile(file2, dtype=np.float16)
    10.     print("output: ", output)
    11.     print("-------------golden is :")
    12.     print("golden: ", golden)
    13.  
    14.     different_element_results = np.isclose(
    15.         output, golden,
    16.         rtol=5e-2,
    17.         atol=1e-3,
    18.         equal_nan=True)
    19.     different_element_indexes = np.where(
    20.         different_element_results != np.array((True,)))[0]
    21.     return 0 if different_element_indexes.size == 0 else 1
    22.  
    23. if __name__ == '__main__':
    24.     intput_file1 = sys.argv[1]
    25.     output_file = sys.argv[2]
    26.     cmp_result = data_compare(intput_file1, output_file)
    27.  
    28.     if (cmp_result == 0):
    29.         sys.exit(0)
    30.     else:
    31.         sys.exit(1)
     

    main.cpp中，需要将CreateOpDesc()函数根据具体的输入输出来做修改。
     

    1. OperatorDesc CreateOpDesc()
    2. {
    3.     std::vector<int64_t> shape{8, 2048};
    4.     aclDataType dataType = ACL_FLOAT16;
    5.     aclFormat format = ACL_FORMAT_ND;
    6.     OperatorDesc opDesc;
    7.     opDesc.AddInputTensorDesc(dataType, shape.size(), shape.data(), format);
    8.     opDesc.AddOutputTensorDesc(dataType, shape.size(), shape.data(), format);
    9.     return opDesc;
    10. }
     

    op_runner.cpp中将两段式API修改为自己算子的API，请善用Ctrl + F搜索关键代码进行修改，具体的API名称可以查看算子目录下的build_out/autogen目录。
     

    1. ...
    2. auto ret = aclnnSinhCustomGetWorkspaceSize(inputTensor_[0], outputTensor_[0], &workspaceSize, &handle);
    3. ...
    4. INFO_LOG("Execute aclnnSinhCustomGetWorkspaceSize success, workspace size %lu", workspaceSize);
    5. ...
    6. if (aclnnSinhCustom(workspace, workspaceSize, handle, stream) != ACL_SUCCESS)
    7. {
    8.     ...
    9. }
    10. INFO_LOG("Execute aclnnSinhCustom success");
    11. ...
     

    接着修改src/CMakeLists.txt。
     

    1. set(AUTO_GEN_PATH "../SinhCustom/build_out/autogen") # 16行
    2.  
    3. # 50行以后，修改可执行文件的名称
    4. add_executable(execute_sinh_op
    5.     ${AUTO_GEN_PATH}/aclnn_sinh_custom.cpp
    6.     operator_desc.cpp
    7.     op_runner.cpp
    8.     main.cpp
    9.     op_runner.cpp
    10.     common.cpp
    11. )
    12.  
    13. target_link_libraries(execute_sinh_op
    14.     ascendcl
    15.     acl_op_compiler
    16.     nnopbase
    17.     stdc++
    18. )
    19.  
    20. install(TARGETS execute_sinh_op DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})
     

    最后修改run.sh脚本中关于路径的部分。
     

    修改完成后，就可以执行run.sh脚本进行单算子API调用了。
     

    1. INFO: acl executable run success!
    2. input1:  [ 0.468  -0.2585 -3.066  ...  0.9136 -1.117  -1.368 ]
    3. output:  [  0.485   -0.2615 -10.71   ...   1.047   -1.365   -1.837 ]
    4. -------------golden is :
    5. golden:  [  0.4854  -0.2615 -10.71   ...   1.046   -1.364   -1.837 ]
    6. INFO: compare golden data success!
     

    出现上述提示证明算子通过验证。
     

     

    5 Ascend C学习资源
     

    Ascend C配套丰富的学习资料，包括教程文档、交流社区、案例代码等，这些资源将帮助您理解Ascend C编程语言的各种概念和技巧，为您的自主学习提供便利。
     

                
    
                    
    
                
    • 

                
  • 

                    相关阅读:
    
                    
asp.net高校网上评教信息系统VS开发sqlserver数据库web结构c#编程计算机网页项目                            
    
网络钓鱼攻击飙升，265个品牌在2022年上半年被冒充                            
    
【技术指南资料】编码器与正交译码器                            
    
保姆级vue-pdf的使用过程                            
    
利用Quartz设计采集系统并实现系统双活机制_在SpringCloud中自己设计系统双活---SpringCloud工作笔记178                            
    
JS字符串知识点                            
    
【Maven】使用maven profile 动态激活不同环境、依赖打包部署                            
    
养老院一键报警的重要性和应用                            
    
基于spingboot通过spark-sql进行大批量插入clickhous分布式数据库                            
    
Android AssetManager初探                            
    
                        
                
    • 
                
  • 
                    原文地址：https://blog.csdn.net/m0_71340392/article/details/134451588
                
    •