运行验证算子工程

环境准备

• 使用Kernel Launch算子工程之前，需要参考环境准备章节安装驱动固件和CANN软件包，完成开发环境和运行环境的准备。
• 使用该算子工程要求cmake版本为3.16及以上版本，如不符合要求，请参考如下的命令示例更新cmake版本，如下示例以更新到3.16.0版本为例。

  wget https://cmake.org/files/v3.16/cmake-3.16.0.tar.gz --no-check-certificate 
  tar -zxvf cmake-3.16.0.tar.gz 
  cd cmake-3.16.0
   ./bootstrap --prefix=/usr 
  sudo make 
  sudo make install

工程目录

您可以单击矢量算子样例，获取核函数开发和运行验证的完整样例。样例目录结构如下所示：

AddKernelInvocationNeo
|-- cmake                                                 // CMake编译文件
|-- scripts
|  ├── gen_data.py                                     // 输入数据和真值数据生成脚本文件
|  ├── verify_result.py                                // 验证输出数据和真值数据是否一致的验证脚本
|-- CMakeLists.txt                                        // CMake编译配置文件
|-- add_custom.cpp                                        // 矢量算子kernel实现
|-- data_utils.h                                          // 数据读入写出函数
|-- main.cpp                                              // 主函数，调用算子的应用程序，含CPU域及NPU域调用
|-- run.sh                                                // 编译运行算子的脚本

基于该算子工程，开发者进行算子开发的步骤如下：

• 完成算子kernel侧实现。
• 编写算子调用应用程序main.cpp。

• 编写CMake编译配置文件CMakeLists.txt。

• 根据实际需要修改输入数据和真值数据生成脚本文件gen_data.py；验证输出数据和真值数据是否一致的验证脚本verify_result.py。
• 根据实际需要修改编译运行算子的脚本run.sh并执行该脚本，完成算子的编译运行和结果验证。

算子kernel侧实现

请参考矢量编程和工程目录中的矩阵算子、融合算子的kernel实现完成Ascend C算子实现文件的编写。

算子调用应用程序

下面代码以固定shape的add_custom算子为例，介绍算子核函数调用的应用程序main.cpp如何编写。您在实现自己的应用程序时，需要关注由于算子核函数不同带来的修改，包括算子核函数名，入参出参的不同等，合理安排相应的内存分配、内存拷贝和文件读写等，相关API的调用方式直接复用即可。

1. 按需包含头文件，通过ASCENDC_CPU_DEBUG宏区分CPU/NPU侧需要包含的头文件。需要注意的是，NPU侧需要包含对应的核函数调用接口声明所在的头文件alcrtlaunch_{kernel_name}.h（该头文件为工程框架自动生成），kernel_name为算子核函数的名称。

   #include "data_utils.h"
   #ifndef ASCENDC_CPU_DEBUG
   #include "acl/acl.h"
   #include "aclrtlaunch_add_custom.h"
   #else
   #include "tikicpulib.h"
   extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z);
   #endif

2. 应用程序框架编写。该应用程序通过ASCENDC_CPU_DEBUG宏区分代码逻辑运行于CPU侧还是NPU侧。

   int32_t main(int32_t argc, char* argv[])
   {
       uint32_t blockDim = 8;
       size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
       size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);
   
   #ifdef ASCENDC_CPU_DEBUG
       // 用于CPU调试的调用程序
   #else
       // NPU侧运行算子的调用程序
   #endif
       return 0;
   }

3. CPU侧运行验证。完成算子核函数CPU侧运行验证的步骤如下：
   图1 CPU侧运行验证步骤
   

       // 使用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);
       // 矢量算子需要设置内核模式为AIV模式
       AscendC::SetKernelMode(KernelMode::AIV_MODE);
       // 调用ICPU_RUN_KF调测宏，完成核函数CPU侧的调用
       ICPU_RUN_KF(add_custom, blockDim, x, y, z);
       // 输出数据写出
       WriteFile("./output/output_z.bin", z, outputByteSize);
       // 调用GmFree释放申请的资源
       AscendC::GmFree((void *)x);
       AscendC::GmFree((void *)y);
       AscendC::GmFree((void *)z);

4. NPU侧运行验证。完成算子核函数NPU侧运行验证的步骤如下：
   图2 NPU侧运行验证步骤
   

       // AscendCL初始化
       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);
       // 将数据从Host上拷贝到Device上
       CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
       CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
       // 用ACLRT_LAUNCH_KERNEL接口调用核函数完成指定的运算
       ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, xDevice, yDevice, zDevice);
       // 用内核调用符<<<>>>调用核函数完成指定的运算,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));
       // AscendCL去初始化
       CHECK_ACL(aclrtDestroyStream(stream));
       CHECK_ACL(aclrtResetDevice(deviceId));
       CHECK_ACL(aclFinalize());

CMake编译配置文件编写

本节会介绍CMake文件中一些关键环境变量和Cmake命令参数的说明，通常情况下不需要开发者修改，但是这些参数可以帮助开发者更好的理解编译原理，方便有能力的开发者对Cmake进行定制化处理。

输入数据和真值数据生成以及验证脚本文件

以固定shape的add_custom算子为例，输入数据和真值数据生成的脚本样例如下：根据算子的输入输出编写脚本，生成输入数据和真值数据。

#!/usr/bin/python3
# -*- coding:utf-8 -*-
# Copyright 2022-2023 Huawei Technologies Co., Ltd
import numpy as np

def gen_golden_data_simple():
    input_x = np.random.uniform(1, 100, [8, 2048]).astype(np.float16)
    input_y = np.random.uniform(1, 100, [8, 2048]).astype(np.float16)
    golden = (input_x + input_y).astype(np.float16)

    input_x.tofile("./input/input_x.bin")
    input_y.tofile("./input/input_y.bin")
    golden.tofile("./output/golden.bin")

if __name__ == "__main__":
    gen_golden_data_simple()

验证输出数据和真值数据是否一致的验证脚本样例如下：当前使用numpy接口计算了输出数据和真值数据的绝对误差和相对误差，误差在容忍偏差范围内，视为精度符合要求，输出"test pass"字样。

import os
import sys
import numpy as np

loss = 1e-3 # 容忍偏差，一般fp16要求绝对误差和相对误差均不超过千分之一
minimum = 10e-10

def verify_result(real_result, golden):
    real_result = np.fromfile(real_result, dtype=np.float16) # 从bin文件读取实际运算结果
    golden = np.fromfile(golden, dtype=np.float16) # 从bin文件读取预期运算结果
    result = np.abs(real_result - golden) # 计算运算结果和预期结果偏差
    deno = np.maximum(np.abs(real_result), np.abs(golden)) # 获取最大值并组成新数组
    result_atol = np.less_equal(result, loss) # 计算绝对误差
    result_rtol = np.less_equal(result / np.add(deno, minimum), loss) # 计算相对误差
    if not result_rtol.all() and not result_atol.all():
        if np.sum(result_rtol == False) > real_result.size * loss and np.sum(result_atol == False) > real_result.size * loss:
            print("[ERROR] result error")
            return False
    print("test pass")
    return True

if __name__ == '__main__':
    verify_result(sys.argv[1],sys.argv[2])

修改并执行一键式编译运行脚本

您可以基于样例工程中提供的一键式编译运行脚本进行快速编译，并在CPU侧和NPU侧执行Ascend C算子。一键式编译运行脚本主要完成以下功能：

图4 一键式编译运行脚本流程图

样例中提供的一键式编译运行脚本并不能适用于所有的算子运行验证场景，使用时请根据实际情况进行修改。

• 根据Ascend C算子的算法原理的不同，自行实现输入和真值数据的生成脚本。