基于样例工程完成核函数（Kernel）直调
说明
本章节介绍的基于样例工程完成核函数（Kernel）直调的方式，后续不再演进。推荐开发者直接使用命令行或者编写Cmake文件进行编译，详细内容请参考AI Core SIMD编译。
下文将以Add矢量算子为例对核函数（Kernel）直调算子开发流程进行详细介绍。
环境准备
使用Kernel Launch算子工程之前，需要参考环境准备章节安装驱动固件和CANN软件包，完成开发环境和运行环境的准备。
使用该算子工程要求cmake版本为3.16及以上版本，如不符合要求，请参考如下的命令示例更新cmake版本，如下示例以更新到3.16.0版本为例。
Textwget https://cmake.org/files/v3.16/cmake-3.16.0.tar.gz 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算子实现文件的编写。
说明
一个算子核函数（Kernel）实现文件中只支持定义一个核函数（Kernel）。
算子调用应用程序
下面代码以固定shape的add_custom算子为例，介绍算子核函数（Kernel）调用的应用程序main.cpp如何编写。您在实现自己的应用程序时，需要关注由于算子核函数（Kernel）不同带来的修改，包括算子核函数（Kernel）名，入参出参的不同等，合理安排相应的内存分配、内存拷贝和文件读写等，相关API的调用方式直接复用即可。
按需包含头文件，通过ASCENDC_CPU_DEBUG宏区分CPU/NPU侧需要包含的头文件。需要注意的是，NPU侧需要包含对应的核函数（Kernel）调用接口声明所在的头文件aclrtlaunch_{kernel_name}.h（该头文件为工程框架自动生成），kernel_name为算子核函数（Kernel）的名称。
Text#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应用程序框架编写。该应用程序通过ASCENDC_CPU_DEBUG宏区分代码逻辑运行于CPU侧还是NPU侧。
Textint32_t main(int32_t argc, char* argv[]) { uint32_t numBlocks = 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; }CPU侧运行验证。完成算子核函数（Kernel）CPU侧运行验证的步骤如下：
图1 CPU侧运行验证步骤
Text// 使用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调测宏，完成核函数（Kernel）CPU侧的调用 ICPU_RUN_KF(add_custom, numBlocks, x, y, z); // 输出数据写出 WriteFile("./output/output_z.bin", z, outputByteSize); // 调用GmFree释放申请的资源 AscendC::GmFree((void *)x); AscendC::GmFree((void *)y); AscendC::GmFree((void *)z);NPU侧运行验证。完成算子核函数（Kernel）NPU侧运行验证的步骤如下：
图2 NPU侧运行验证步骤
Text// 初始化 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)); // 用内核调用符<<<>>>调用核函数（Kernel）完成指定的运算,add_custom_do中封装了<<<>>>调用 add_custom_do(numBlocks, nullptr, stream, xDevice, yDevice, zDevice); // 用ACLRT_LAUNCH_KERNEL接口调用核函数（Kernel）完成指定的运算 // ACLRT_LAUNCH_KERNEL(add_custom)(numBlocks, 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());说明 针对<<<>>>调用方式在核函数（Kernel）章节已有说明，这里仅对ACLRT_LAUNCH_KERNEL调用接口的使用方法介绍如下：
TextACLRT_LAUNCH_KERNEL(kernel_name)(numBlocks, stream, argument list);- kernel_name：算子核函数（Kernel）的名称。
- numBlocks：规定了核函数（Kernel）将会在几个核上执行。每个执行该核函数（Kernel）的核会被分配一个逻辑ID，即block_idx，可以在核函数（Kernel）的实现中调用GetBlockIdx来获取block_idx。
- stream，类型为aclrtStream，stream用于维护一些异步操作的执行顺序，确保按照应用程序中的代码调用顺序在Device上执行。stream创建等管理接口请参考《Runtime运行时API》。
- argument list：参数列表，与核函数（Kernel）的参数列表保持一致。
CMake编译配置文件编写
本节会介绍CMake文件中一些关键环境变量和Cmake命令参数的说明，通常情况下不需要开发者修改，但是这些参数可以帮助开发者更好的理解编译原理，方便有能力的开发者对Cmake进行定制化处理。
表1 环境变量说明
环境变量 | 配置说明 |
|---|---|
SOC_VERSION | AI处理器的型号。 |
ASCEND_CANN_PACKAGE_PATH | CANN软件包安装后的实际路径。 |
CMAKE_BUILD_TYPE | 编译模式选项，可配置为：
|
CMAKE_INSTALL_PREFIX | 用于指定CMAKE执行install时，安装的路径前缀，执行install后编译产物（ascendc_library中指定的target以及对应的头文件）会安装在该路径下。默认路径为当前目录的out目录下。 |
CMAKE_CXX_COMPILER_LAUNCHER | 用于配置C++语言编译器（如g++）、毕昇编译器的启动器程序为ccache，配置后即可开启cache缓存编译，加速重复编译并提高构建效率。使用该功能前需要安装ccache。 配置方法如下，在对应的CMakeLists.txt进行设置： set(CMAKE_CXX_COMPILER_LAUNCHER <launcher_program>) 其中<launcher_program>是ccache的安装路径，比如ccache的安装路径为/usr/bin/ccache，示例如下： set(CMAKE_CXX_COMPILER_LAUNCHER /usr/bin/ccache) |
表2 Cmake命令语法说明
Cmake命令 | 语法说明 |
|---|---|
add_executable | 使用指定的源文件将可执行文件添加到项目中。和Cmake通用的命令参数使用方法一致。 |
ascendc_library | 使用指定的核函数（Kernel）源文件向项目（project）添加库。语法格式如下： ascendc_library(<target_name> [STATIC | SHARED]
[<source>...]) 其中<target_name>表示库文件的名字，该库文件会根据命令里列出的源文件来建立。STATIC、SHARED的作用是指定生成的库文件的类型。STATIC库是目标文件的归档文件，在连接其它目标的时候使用。SHARED库会被动态连接（动态连接库），在运行时会被加载。<source>表示核函数（Kernel）源文件。 |
ascendc_fatbin_library | 使用指定的核函数（Kernel）源文件编译生成一个核函数（Kernel）二进制文件，供核函数（Kernel）加载和执行接口使用。语法格式如下： ascendc_fatbin_library(<target_name> [<source>...])
说明：
|
ascendc_compile_definitions | 添加编译宏。可以添加Ascend C提供的编译宏和开发者自定义的编译宏。语法格式如下： ascendc_compile_definitions(<target_name> [PRIVATE]
[<xxx>...]) Ascend C提供的编译宏介绍如下：
|
ascendc_compile_options | 添加编译选项。可以添加相应的编译选项用于host侧与device侧的编译过程。语法格式如下： ascendc_compile_options(<target_name> PRIVATE
[<xxx>...]
)默认情况下，指定的编译选项都将传递给device侧编译器进行编译。若想传递编译选项给host侧编译器，则需要使用“-forward-options-to-host-compiler”编译选项，该选项后的编译选项将传递给host侧编译器，示例如下： ascendc_compile_options(<target_name> PRIVATE
-g
-forward-options-to-host-compiler
-gdwarf-4
)如上代码所示，在编译时，“-g”编译选项传递给device侧编译器，“-gdwarf-4”编译选项传递给host侧编译器。 备注：host侧编译选项只支持g++与clang编译器共同支持的编译选项。 |
ascendc_include_directories | 添加开发者自定义的头文件搜索路径。语法格式如下： ascendc_include_directories(<target_name> [PRIVATE]
[<xxx>...]) |
简化的编译流程图如下图所示：将算子核函数（Kernel）源文件编译生成kernel侧的库文件（*.so或*.a库文件）；工程框架自动生成核函数（Kernel）调用接口声明头文件；编译main.cpp（算子调用应用程序）时依赖上述头文件，将编译应用程序生成的目标文件和kernel侧的库文件进行链接，生成最终的可执行文件。
图3 编译简化流程图
编译安装结束后在CMAKE_INSTALL_PREFIX目录下生成的编译产物示例如下；最终的可执行文件会生成在cmake命令的执行目录下。
out
├── lib
│ ├── libkernels1.a
│ ├── libkernels2.so
├── include
│ ├── kernels1
│ ├── aclrtlaunch_matmul_custom.h
│ ├── aclrtlaunch_add_custom.h
│ ├── kernels2
│ ├── aclrtlaunch_xxx.h
│ ├── ...
对于lib目录下生成的库文件可通过msobjdump工具进一步解析得到kernel信息，具体操作参见msobjdump工具。
输入数据和真值数据生成以及验证脚本文件
以固定shape的add_custom算子为例，输入数据和真值数据生成的脚本样例如下：根据算子的输入输出编写脚本，生成输入数据和真值数据。
#!/usr/bin/python3
# -*- coding:utf-8 -*-
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算子的算法原理的不同，自行实现输入和真值数据的生成脚本。
完成上述文件的编写后，可以执行一键式编译运行脚本，编译和运行应用程序。
脚本执行方式和脚本参数介绍如下：
bash run.sh --run-mode=npu --soc-version=<soc_version> --install-path=<install_path> --build-type=Debug --install-prefix=<install-prefix>
bash run.sh -r npu -v <soc_version> -i <install_path> -b Debug -p <install-prefix>
表3 脚本参数介绍
参数名 | 参数简写 | 参数介绍 |
|---|---|---|
--run-mode | -r | 表明算子以cpu模式或npu模式运行。 取值为cpu或npu。 |
--soc-version | -v | 算子运行的AI处理器型号。 |
--install-path | -i | 配置为CANN软件的安装路径，请根据实际安装路径进行修改。 默认值为$HOME/Ascend/ascend-toolkit/latest。 |
--build-type | -b | 编译模式选项，可配置为：
默认值为Debug。 |
--install-prefix | -p | 用于指定CMAKE执行install时，安装的路径前缀，执行install后编译产物（ascendc_library中指定的target以及对应的头文件）会安装在该路径下。默认路径为当前目录的out目录下。 |
脚本执行完毕输出"test pass"字样表示算子精度符合要求。
CPU侧验证核函数（Kernel）
在非昇腾设备上，开发者可以利用CPU仿真环境先行进行算子开发和测试，并在准备就绪后，利用昇腾设备进行加速计算。在编译与运行章节，我们已经介绍了算子核函数（Kernel）程序NPU域的编译运行。类似地，CPU域调试也是通过毕昇编译器编译算子核函数（Kernel）程序。此时算子核函数（Kernel）程序链接CPU调测库，执行编译生成的可执行文件，可以完成算子CPU域的运行验证。CPU侧的运行程序，通过GDB通用调试工具进行单步调试，可以精准验证程序执行流程是否符合预期。
图5 CPU域和NPU域的核函数（Kernel）运行逻辑对比
基于核函数（Kernel）直调样例工程，通过ACLRT_LAUNCH_KERNEL接口调用核函数（Kernel）时，可实现CPU与NPU域的代码的统一，且该方式仅支持以下型号：
- Atlas A3 训练系列产品/Atlas A3 推理系列产品
- Atlas A2 训练系列产品/Atlas A2 推理系列产品
- Atlas 推理系列产品
下面代码以add_custom算子为例，介绍算子核函数（Kernel）在CPU侧验证时，算子调用的应用程序如何编写（通过ACLRT_LAUNCH_KERNEL接口调用核函数（Kernel）的方式）。您在实现自己的应用程序时，需要关注由于算子核函数（Kernel）不同带来的修改，包括算子核函数（Kernel）名，入参出参的不同等，合理安排相应的内存分配、内存拷贝和文件读写等，相关API的调用方式直接复用即可。
按需包含头文件。
Text#include "data_utils.h" #include "acl/acl.h" #include "aclrtlaunch_add_custom.h"应用程序框架编写。
Textint32_t main(int32_t argc, char* argv[]) { uint32_t numBlocks = 8; size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); // 运行算子的调用程序 return 0; }运行验证。
Text// 初始化 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接口调用核函数（Kernel）完成指定的运算 ACLRT_LAUNCH_KERNEL(add_custom)(numBlocks, 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());
说明 为了实现CPU域与NPU域代码归一，仅对部分acl接口进行适配，开发者在使用CPU域调测功能时，仅支持使用如下acl接口：
- 有实际功能接口，支持CPU域调用
- aclDataTypeSize、aclFloat16ToFloat、aclFloatToFloat16。
- aclrtMalloc、aclrtFree、aclrtMallocHost、aclrtFreeHost、aclrtMemset、aclrtMemsetAsync、aclrtMemcpy、aclrtMemcpyAsync、aclrtMemcpy2d、aclrtMemcpy2dAsync、aclrtCreateContext、aclrtDestroyContext。
- 无实际功能接口，打桩实现。
- Profiling数据采集 aclprofInit、aclprofSetConfig、aclprofStart、aclprofStop、aclprofFinalize。
- 系统配置 aclInit、aclFinalize、aclrtGetVersion。
- 运行时管理 aclrtSetDevice、aclrtResetDevice、aclrtCreateStream、aclrtCreateStreamWithConfig、aclrtDestroyStream、aclrtDestroyStreamForce、aclrtSynchronizeStream、aclrtCreateContext、aclrtDestroyContext。