本节介绍的核函数运行验证功能，主要目的是帮助开发者快速的理解矢量编程的编程模型、熟悉矢量算子的开发和基础调用流程。所以本节仅提供简单的算子运行验证功能，不支持算子的shape、数据类型是动态的场景，也不支持获取用户的workspace特性。如果您想要使用上述动态shape等特性，请参考算子开发（进阶篇）进行进一步的学习。

核函数即算子kernel程序开发完成后，即可编写host侧的核函数调用程序，实现从host侧的APP程序调用算子，进行运行验证。本节将会介绍CPU侧和NPU侧两种运行验证方法：

CPU侧运行验证主要通过ICPU_RUN_KF CPU调测宏等CPU调测库提供的接口来完成；
NPU侧运行验证主要通过使用<<<>>>内核调用符和AscendCL API提供的运行时接口来完成。

CPU侧和NPU侧的运行验证原理图如下：

图1 CPU侧和NPU侧运行验证原理图

您可以根据下文的介绍来完成基本的运行验证流程，流程中使用到接口的详细细节请参考：

CPU调测库请参考调测接口。
<<<>>>内核调用符使用方法请参考核函数调用。
AscendCL API使用方法请参考《应用软件开发指南（C&C++）》“AscendCL API参考”章节。

基于NPU域算子的调用接口（<<<>>>内核调用符）编写的算子程序，通过毕昇编译器编译后运行，可以完成算子NPU域的运行验证；基于CPU域算子的调用接口（ICPU_RUN_KF CPU）编写的算子程序，通过标准的GCC编译器进行编译后运行，可以完成算子CPU域的运行验证。

CPU侧的运行程序，通过GDB通用调试工具进行单步调试，精准验证程序执行流程是否符合预期。如果您想进一步了解CPU侧调试的具体内容，可在完成本节内容的学习后参考CPU域调试。

创建代码目录

您可以单击LINK，获取核函数开发和运行验证的完整样例。

代码目录如下：

Add
|-- input                                                 // 存放脚本生成的输入数据目录
|-- output                                                // 存放算子运行输出数据和真值数据的目录
|-- CMakeLists.txt -> ../kernel_template/CMakeLists.txt   // 编译工程文件
|-- add_custom.cpp                                      // 算子kernel实现
|-- add_custom.py                                       // 输入数据和真值数据生成脚本文件
|-- cmake -> ../kernel_template/cmake                     // 编译工程文件
|-- data_utils.h -> ../kernel_template/data_utils.h       // 数据读入写出函数
|-- main.cpp                                             // 主函数，调用算子的应用程序，含CPU域及NPU域调用
|-- run.sh -> ../kernel_template/run.sh                  // 编译运行算子的脚本

在进行算子调用前，请确保已经参考矢量编程完成了Ascend C算子实现文件的编写，本样例中为add_custom.cpp文件。除此之外，还需要特别关注以下文件，需要根据自己实际的使用场景进行修改。

调用算子的应用程序：main.cpp。
输入数据和真值数据生成脚本文件：add_custom.py。
编译运行算子的脚本：run.sh。

算子调用应用程序

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

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

#include "data_utils.h"
#ifndef __CCE_KT_TEST__
#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 __CCE_KT_TEST__
    // 用于CPU调试的调用程序
#else
    // NPU侧运行算子的调用程序
#endif
    return 0;
}

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

图2 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);

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

图3 NPU侧运行验证步骤

    // AscendCL初始化
    CHECK_ACL(aclInit(nullptr));
    // 运行管理资源申请
    aclrtContext context;
    int32_t deviceId = 0;
    CHECK_ACL(aclrtSetDevice(deviceId));
    CHECK_ACL(aclrtCreateContext(&context, 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));
    // AscendCL去初始化
    CHECK_ACL(aclrtDestroyStream(stream));
    CHECK_ACL(aclrtDestroyContext(context));
    CHECK_ACL(aclrtResetDevice(deviceId));
    CHECK_ACL(aclFinalize());

输入数据和真值数据生成脚本文件

根据算子的输入输出编写脚本，生成输入数据和真值数据。为了方便后续的调用，该脚本文件名和算子实现文件名相同，比如算子实现文件名为KERNEL_NAME.cpp，脚本文件应命名为KERNEL_NAME.py。本样例中脚本文件名为add_custom.py。

此处以固定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(-100, 100, [8, 2048]).astype(np.float16)
    input_y = np.random.uniform(-100, 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()

一键式编译运行脚本

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

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

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

run.sh中ASCEND_HOME_DIR为CANN软件的安装位置，请替换为实际安装路径。例如，$HOME/Ascend/ascend-toolkit/latest。
根据Ascend C算子的算法原理的不同，自行实现输入和真值数据的生成脚本。
由于每个算子的输入输出个数各不相同，需修改run.sh中数据对比部分，当前使用md5sum对比了所有输出bin文件。此处使用md5sum进行结果比对的方式仅作为示例，用户可自己针对使用场景选择合适的精度比对方式，比如使用numpy中的接口自行实现。

运行验证

完成上述文件的编写后，可以执行一键式编译运行脚本，编译和运行应用程序。脚本执行方式和脚本参数介绍如下：

bash run.sh <kernel_name> <soc_version> <core_type> <run_mode>

表1 脚本参数介绍
参数名	参数介绍	取值
<kernel_name>	Ascend C算子实现文件的文件名	比如Add算子实现文件为add_custom.cpp，则应传入add_custom。
*<soc_version>*	算子运行的AI处理器型号	根据实际算子运行的AI处理器型号进行配置。
<core_type>	表明算子在AiCore上或者VectorCore上运行	AiCore或VectorCore
<run_mode>	表明算子以cpu模式或npu模式运行	cpu或npu

如下图所示，脚本执行完毕会出现如下示例的md5值打印，实际输出结果和真值数据的md5值一致表示算子精度符合要求。

INFO：execute op on cpu succeed!
md5sum:
6a99e41a84b14dd04f32730ceb9a3988  output/golden.bin
6a99e41a84b14dd04f32730ceb9a3988  output/output_y.bin