烟台赶集网网站建设,养老网站建设合同,android下载安装官方免费下载,深圳服装设计公司排名前十强本文分享自《AscendC算子开发及单算子调用》#xff0c;作者#xff1a;goldpancake。
笔者在阅读Ascend C官方文档的过程中发现#xff0c;对于初学者来说#xff0c;尤其是第一次接触异构编程思想的初学者#xff0c;有部分内容是无需特别关注的#xff0c;例如算子工…本文分享自《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环境中的解释器。
#!/root/miniconda3/bin/python3
# codingutf-8Function:
This file mainly involves main function of op generation module.
Copyright Information:
Huawei Technologies Co., Ltd. All Rights Reserved © 2020原来是root用户下的conda配置为了默认激活base环境笔者安装时没有注意这一点导致在CANN包安装的过程中选择到了conda环境下的python解释器这样一来其他用户肯定是没有权限的。在关闭base环境重新安装CANN包后问题解决。 2 算子开发流程
至此环境准备好后开始正式的算子开发步骤。 2.1 算子工程配置文件
CANN包中提供了一个自动生成算子工程的工具msopgen该工具可以通过一个json配置文件来生成完整的算子工程具体的编写方式请参考Ascend C官方文档。
这里以sinh算子为例该算子是一元操作所以只需要一个输入且输出形状与输入形状一致。根据该特征来编写json文件为了贴合Ascend C官方建议的编程范式将文件命名为sinh_custom.json。为了简洁这里我们只实现一种数据类型的操作。
[{op: SinhCustom,language: cpp,input_desc: [{name: x,param_type: required,format: [ND],type: [fp16]}],output_desc: [{name: y,param_type: required,format: [ND],type: [fp16]}]}
] 2.2 生成算子工程
创建一个文件夹用作算子工程目录使用msopgen工具执行如下命令来生成算子工程。
mkdir /path/to/SinhCustom
/path/to/msopgen gen -i /path/to/sinh_custom.json -c ai_core-Ascend910 -lan cpp -out /path/to/SinhCustom
命令行会输出类似如下的信息
2023-10-07 14:58:42 (942445) - [INFO] Start to generate AI Core operator files.
2023-10-07 14:58:42 (942445) - [INFO] Start to parse the ir template:/path/to/SinhCustom/sinh_custom.json
2023-10-07 14:58:42 (942445) - [INFO] Start to parse the op: SinhCustom
2023-10-07 14:58:42 (942445) - [INFO] Start to parse the input_desc: x
2023-10-07 14:58:42 (942445) - [INFO] Start to parse the output_desc: y
2023-10-07 14:58:42 (942445) - [WARNING] The attr value is invalid or no attr exists in the map.
2023-10-07 14:58:42 (942445) - [INFO] Start to check the type and format between the inputs/outputs in IR template.
2023-10-07 14:58:42 (942445) - [INFO] Start to generate a new project.
2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/cmake/config.cmake generated successfully.
2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/op_host/sinh_custom_tiling.h generated successfully.
2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/op_host/sinh_custom.cpp generated successfully.
2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/op_kernel/sinh_custom.cpp generated successfully.
2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/framework/tf_plugin/tensorflow_sinh_custom_plugin.cc generated successfully.
2023-10-07 14:58:42 (942445) - [INFO] File /path/to/SinhCustom/framework/tf_plugin/CMakeLists.txt generated successfully.
2023-10-07 14:58:42 (942445) - [INFO] Generation completed.
此时会发现指定的输出目录只已经生成了一系列的算子工程文件。
SinhCustom
├── build.sh
├── cmake
├── CMakeLists.txt
├── CMakePresets.json # 这个配置项需要修改
├── framework
├── op_host
│ ├── CMakeLists.txt
│ ├── sinh_custom.cpp # 算子host侧核心逻辑
│ └── sinh_custom_tiling.h # 算子tiling结构体定义
├── op_kernel
│ ├── CMakeLists.txt
│ └── sinh_custom.cpp # 算子kernel侧核心逻辑
├── scripts
└── 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代码如下
#ifndef SINH_CUSTOM_TILING_H // 头文件保护记得加上自动生成的文件中不包含
#define SINH_CUSTOM_TILING_H
#include register/tilingdata_base.hnamespace optiling
{BEGIN_TILING_DATA_DEF(TilingData)TILING_DATA_FIELD_DEF(uint32_t, formerNum); // 整块个数TILING_DATA_FIELD_DEF(uint32_t, tailNum); // 尾块个数TILING_DATA_FIELD_DEF(uint32_t, formerLength); // 整块内元素个数TILING_DATA_FIELD_DEF(uint32_t, tailLength); // 尾块内元素个数TILING_DATA_FIELD_DEF(uint32_t, alignNum); // 最小对齐单位元素个数END_TILING_DATA_DEF;REGISTER_TILING_DATA_CLASS(SinhCustom, TilingData)
}#endif 然后在op_host/sinh_custom.cpp中实现具体的tiling策略代码如下 namespace optiling
{constexpr uint32_t BLOCK_DIM 24; // 划分核心数量constexpr uint32_t SIZE_OF_HALF 2; // 数据类型的字节数constexpr uint32_t BLOCK_SIZE 32; // 昇腾设备上的数据block为32字节constexpr uint32_t ALIGN_NUM BLOCK_SIZE / SIZE_OF_HALF; // 最小对齐单位static ge::graphStatus TilingFunc(gert::TilingContext *context){TilingData tiling;uint32_t totalLength context-GetInputTensor(0)-GetShapeSize();context-SetBlockDim(BLOCK_DIM);// 使输入向上对齐uint32_t totalLengthAligned ((totalLength ALIGN_NUM - 1) / ALIGN_NUM) * ALIGN_NUM;// 计算整块和尾块个数uint32_t formerNum (totalLengthAligned / ALIGN_NUM) % BLOCK_DIM;uint32_t tailNum BLOCK_DIM - formerNum;// 计算整块和尾块的元素个数uint32_t formerLength ((totalLengthAligned / BLOCK_DIM ALIGN_NUM - 1) / ALIGN_NUM) * ALIGN_NUM;uint32_t tailLength (totalLengthAligned / BLOCK_DIM / ALIGN_NUM) * ALIGN_NUM;// 设置tiling参数tiling.set_formerNum(formerNum);tiling.set_tailNum(tailNum);tiling.set_formerLength(formerLength);tiling.set_tailLength(tailLength);tiling.set_alignNum(ALIGN_NUM);// 以下为固定写法不用纠结tiling.SaveToBuffer(context-GetRawTilingData()-GetData(), context-GetRawTilingData()-GetCapacity());context-GetRawTilingData()-SetDataSize(tiling.GetDataSize());context-SetTilingKey(1);size_t *currentWorkspace context-GetWorkspaceSizes(1);currentWorkspace[0] 0;return ge::GRAPH_SUCCESS;}
}
2.3.2 kernel侧实现
有了上述实现的tiling策略我们就可以根据数据划分的逻辑来确定kernel侧的具体实现。根据官方推荐的矢量编程范式我们可以先将算子类的框架写出来再慢慢填充内容。在op_kernel/sinh_custom.cpp中写出算子类框架。
using namespace AscendC; // 记得开启AscendC命名空间
constexpr int32_t BUFFER_NUM 2; // TQue的缓冲数量此处开启双Bufferclass KernelSinh
{
public:__aicore__ inline KernelSinh() {} // 类构造函数无须任何代码__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, // 初始化函数的参数为输入、输出uint32_t formerNum, uint32_t tailNum, // 以及上面定义的一系列tiling参数uint32_t formerLength, uint32_t tailLength,uint32_t alignNum) { /* TODO */ }__aicore__ inline void Process() { /* TODO */ }private:__aicore__ inline void CopyIn() { /* TODO */ }__aicore__ inline void Compute() { /* TODO */ }__aicore__ inline void CopyOut() { /* TODO */ }private:/* TODO */
};
第一步分析算子类的私有数据成员。
首先一定需要的是用来管理内存的Tpipe同时需要输入输出分别对应的TQue和GlobalTensor同时每个逻辑核还需要直到当前处理的数据个数所以需要一个变量tileLength来确定分片大小。
第二步分析算子。
公式$$ {\bf y}\text{sinh}({\bf x})\frac{e^{\bf x}-e^{-{\bf x}}}{2.0} $$
可以观察到我们需要计算两个中间结果分别是$e^{\bf x}$和$e^{-{\bf x}}$所以需要相应的数据结构来存放这两个中间结果Ascend C提供的TBuf可以很好的承担这一责任。
至此我们就将算子类需要的私有数据成员确定了下来。
TPipe pipe; // 用于操作队列
TBufQuePosition::VECCALC tempBuf; // 存放中间结果
TQueQuePosition::VECIN, BUFFER_NUM inQueueX; // 输入队列
TQueQuePosition::VECOUT, BUFFER_NUM outQueueY; // 输出队列
GlobalTensorDTYPE_X xGm; // 输入数据对应的GM内存空间
GlobalTensorDTYPE_Y yGm; // 输出数据对应的GM内存空间
uint32_t tileLength; // 每个逻辑核需要知道分片数据个数
第三步完善算子类的初始化函数Init()。
在该函数中我们需要为GlobalTensor分配内存并初始化相应的TQue同时需要针对某些变量做合法性判断。
__aicore__ inline void Init(GM_ADDR x, GM_ADDR y,uint32_t formerNum, uint32_t tailNum,uint32_t formerLength, uint32_t tailLength,uint32_t alignNum)
{if (GetBlockIdx() formerNum){// 处理整块逻辑this-tileLength formerLength;xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x formerLength * GetBlockIdx(), formerLength);yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y formerLength * GetBlockIdx(), formerLength);}else{// 处理尾块逻辑this-tileLength tailLength;xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x formerLength * formerNum tailLength * (GetBlockIdx() - formerNum), tailLength);yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y formerLength * formerNum tailLength * (GetBlockIdx() - formerNum), tailLength);}ASSERT(alignNum ! 0 align num can not be zero!);pipe.InitBuffer(inQueueX, BUFFER_NUM, (((this-tileLength alignNum - 1) / alignNum) * alignNum) * sizeof(half));pipe.InitBuffer(outQueueY, BUFFER_NUM, (((this-tileLength alignNum - 1) / alignNum) * alignNum) * sizeof(half));
}
第四步完成算子最核心的部分根据矢量编程范式实现算子计算逻辑。
__aicore__ inline void CopyIn()
{LocalTensorDTYPE_X xLocal inQueueX.AllocTensorDTYPE_X();DataCopy(xLocal, xGm, this-tileLength); // GM - LMinQueueX.EnQueDTYPE_X(xLocal);
}
__aicore__ inline void Compute()
{LocalTensorDTYPE_X xLocal inQueueX.DeQueDTYPE_X();LocalTensorDTYPE_Y yLocal outQueueY.AllocTensorDTYPE_Y();pipe.InitBuffer(tempBuf, this-tileLength * sizeof(DTYPE_X));LocalTensorDTYPE_X tempLocal tempBuf.GetDTYPE_X(this-tileLength);// 计算exp(x)Exp(yLocal, xLocal, this-tileLength);// 计算-xhalf nagOne(-1.0);Muls(tempLocal, xLocal, nagOne, this-tileLength);// 计算exp(-x)Exp(tempLocal, tempLocal, this-tileLength);// 计算exp(x)-exp(-x)Sub(yLocal, yLocal, tempLocal, this-tileLength);// 计算最终结果half denominator(0.5);Muls(yLocal, yLocal, denominator, this-tileLength);outQueueY.EnQueDTYPE_Y(yLocal);inQueueX.FreeTensor(xLocal);
}
__aicore__ inline void CopyOut()
{LocalTensorDTYPE_Y yLocal outQueueY.DeQueDTYPE_Y();DataCopy(yGm, yLocal, this-tileLength); // LM - GMoutQueueY.FreeTensor(yLocal);
}
实现的具体细节与接口可以参考Ascend C官方文档。
第五步将Process()函数补全并完善核函数。
__aicore__ inline void Process()
{CopyIn();Compute();CopyOut();
}extern C __global__ __aicore__ void
sinh_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{GET_TILING_DATA(tiling_data, tiling);KernelSinh op;op.Init(x, y,tiling_data.formerNum, tiling_data.tailNum,tiling_data.formerLength, tiling_data.tailLength,tiling_data.alignNum);if (TILING_KEY_IS(1)){op.Process();}
}
至此就完成了kernel侧的实现。 2.3.3 host侧实现
我们回到op_host/sinh_custom.cpp关于类型推导函数这个算子输入输出的形状一致。msopgen生成的算子工程中默认即为输入输出形状一致所以无须改动。如果在写其他复杂算子的时候需要仔细分析数据形状的变化。关于算子原型注册也无须改动。
现在就完成了整个算子的逻辑可以执行build.sh来验证有没有编译时错误若没有错误则可以进行运行时验证。 3 核函数调用
笔者直接将官方的核函数调用样例拿来做了一些修改需要修改的地方如下。
kernel_invocation
├── cmake
├── CMakeLists.txt
├── data_utils.h
├── input
├── main.cpp # 需要修改
├── output
├── run.sh # 需要修改
├── add_custom.cpp # 替换为自己的算子实现
├── add_custom.py # 需要修改
└── verify_result.py # 添加的代码用于验证结果
首先将官方样例中的add_custom.cpp替换为自己实现的kernel侧算子笔者这里的名称为sinh_custom.cpp。同时为了CPU侧调试需要添加一个核函数的包装函数代码如下。
#ifndef __CCE_KT_TEST__
void sinh_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y)
{sinh_customblockDim, l2ctrl, stream(x, y);
}
#endif
注意为了快速验证逻辑在核函数验证过程中未使用动态tiling所以没有之前提到的那些tiling参数。
然后是sinh_custom.py官方样例中是add_custom.py这里修改文件名称因为后面的run.sh中是通过算子文件名来调用这一python脚本的。
由于本算子只需要一个输入向量所以只生成一个input数据然后修改golden数据的生成方式调用numpy中与算子功能相同的函数来计算注意数据类型代码如下。
import numpy as npdef gen_golden_data_simple():np.random.seed(42)input_x np.random.randn(8, 2048).astype(np.float16)golden np.sinh(input_x).astype(np.float16)print(f-----------------------{input_x[0][0]})input_x.tofile(./input/input_x.bin)golden.tofile(./output/golden.bin)if __name__ __main__:gen_golden_data_simple()
main.cpp中要调整相应的内存申请等操作只需要一个inputCPU侧调试和NPU侧调试的代码都需要修改具体如下。
#include stdio.h#include data_utils.h
#ifndef __CCE_KT_TEST__
#include acl/acl.h
extern void sinh_custom_do(uint32_t coreDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y);
#else
#include tikicpulib.h
extern C __global__ __aicore__ void sinh_custom(GM_ADDR x, GM_ADDR y);
#endifint32_t main(int32_t argc, char *argv[])
{size_t inputByteSize 8 * 2048 * sizeof(uint16_t);size_t outputByteSize 8 * 2048 * sizeof(uint16_t);uint32_t blockDim 8;#ifdef __CCE_KT_TEST__uint8_t *x (uint8_t *)AscendC::GmAlloc(inputByteSize);uint8_t *y (uint8_t *)AscendC::GmAlloc(outputByteSize);ReadFile(./input/input_x.bin, inputByteSize, x, inputByteSize);AscendC::SetKernelMode(KernelMode::AIV_MODE);ICPU_RUN_KF(sinh_custom, blockDim, x, y);WriteFile(./output/output_y.bin, y, outputByteSize);AscendC::GmFree((void *)x);AscendC::GmFree((void *)y);
#elseCHECK_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));uint8_t *xHost, *yHost;uint8_t *xDevice, *yDevice;CHECK_ACL(aclrtMallocHost((void **)(xHost), inputByteSize));CHECK_ACL(aclrtMallocHost((void **)(yHost), outputByteSize));CHECK_ACL(aclrtMalloc((void **)xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));CHECK_ACL(aclrtMalloc((void **)yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));ReadFile(./input/input_x.bin, inputByteSize, xHost, inputByteSize);CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));sinh_custom_do(blockDim, nullptr, stream, xDevice, yDevice);CHECK_ACL(aclrtSynchronizeStream(stream));CHECK_ACL(aclrtMemcpy(yHost, outputByteSize, yDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));WriteFile(./output/output_y.bin, yHost, outputByteSize);CHECK_ACL(aclrtFree(xDevice));CHECK_ACL(aclrtFree(yDevice));CHECK_ACL(aclrtFreeHost(xHost));CHECK_ACL(aclrtFreeHost(yHost));CHECK_ACL(aclrtDestroyStream(stream));CHECK_ACL(aclrtDestroyContext(context));CHECK_ACL(aclrtResetDevice(deviceId));CHECK_ACL(aclFinalize());
#endifreturn 0;
} 原样例中的验证方式是求md5和但由于核函数中调用了Exp、Muls等API所以精度可能会有损失不适合用md5sum的方式来验证。这里就需要引入新的文件verify_result.py这里使用了numpy.isclose函数来进行验证这也是官方单算子API调用的结果验证方式。
import sys
import math
import numpy as npdef data_compare(file1, file2,file3):input1 np.fromfile(file1, dtypenp.float16)print(input1: , input1)golden np.fromfile(file2, dtypenp.float16)output np.fromfile(file3, dtypenp.float16)print(output: , output)print(-------------golden is :)print(golden: , golden)different_element_results np.isclose(output, golden,rtol5e-2,atol1e-3,equal_nanTrue)different_element_indexes np.where(different_element_results ! np.array((True,)))[0]if different_element_indexes.size 0:print(result correct!)else:print(result error!)return 0 if different_element_indexes.size 0 else 1if __name__ __main__:intput_file1 sys.argv[1]golden_file sys.argv[2]output_file sys.argv[3]cmp_result data_compare(intput_file1, golden_file, output_file)if (cmp_result 0):sys.exit(0)else: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调用样例拿来做了一些修改。需要修改的几处关键代码如下。
aclnn_online_model
├── build
├── inc
├── README.md
├── run
│ └── out
│ ├── execute_sinh_op
│ ├── result_files
│ └── test_data
│ ├── config
│ └── data
│ ├── generate_data.py # 生成测试数据脚本需要修改
├── run.sh # 需要修改
├── scripts
│ └── verify_result.py # 调整验证方式例如相对和绝对误差参数等
└── src├── CMakeLists.txt # 需要修改├── common.cpp├── main.cpp # 需要修改├── operator_desc.cpp└── op_runner.cpp # 需要修改
具体细节如下。
generate_data.py中按照算子来修改测试数据生成方式。本算子需要half类型的测试数据故代码改为
import numpy as np
a np.random.randn(8, 2048).astype(np.float16)
a.tofile(input_0.bin)
verify_result.py中根据实际读取的输入和输出利用np.isclose来进行比较该函数详细用法参考numpy官方文档。
import sys
import math
import numpy as npdef data_compare(file1, file2):input1 np.fromfile(file1, dtypenp.float16)print(input1: , input1)golden np.sinh(input1).astype(np.float16)output np.fromfile(file2, dtypenp.float16)print(output: , output)print(-------------golden is :)print(golden: , golden)different_element_results np.isclose(output, golden,rtol5e-2,atol1e-3,equal_nanTrue)different_element_indexes np.where(different_element_results ! np.array((True,)))[0]return 0 if different_element_indexes.size 0 else 1if __name__ __main__:intput_file1 sys.argv[1]output_file sys.argv[2]cmp_result data_compare(intput_file1, output_file)if (cmp_result 0):sys.exit(0)else:sys.exit(1)
main.cpp中需要将CreateOpDesc()函数根据具体的输入输出来做修改。
OperatorDesc CreateOpDesc()
{std::vectorint64_t shape{8, 2048};aclDataType dataType ACL_FLOAT16;aclFormat format ACL_FORMAT_ND;OperatorDesc opDesc;opDesc.AddInputTensorDesc(dataType, shape.size(), shape.data(), format);opDesc.AddOutputTensorDesc(dataType, shape.size(), shape.data(), format);return opDesc;
}
op_runner.cpp中将两段式API修改为自己算子的API请善用Ctrl F搜索关键代码进行修改具体的API名称可以查看算子目录下的build_out/autogen目录。
...
auto ret aclnnSinhCustomGetWorkspaceSize(inputTensor_[0], outputTensor_[0], workspaceSize, handle);
...
INFO_LOG(Execute aclnnSinhCustomGetWorkspaceSize success, workspace size %lu, workspaceSize);
...
if (aclnnSinhCustom(workspace, workspaceSize, handle, stream) ! ACL_SUCCESS)
{...
}
INFO_LOG(Execute aclnnSinhCustom success);
...
接着修改src/CMakeLists.txt。
set(AUTO_GEN_PATH ../SinhCustom/build_out/autogen) # 16行# 50行以后修改可执行文件的名称
add_executable(execute_sinh_op${AUTO_GEN_PATH}/aclnn_sinh_custom.cppoperator_desc.cppop_runner.cppmain.cppop_runner.cppcommon.cpp
)target_link_libraries(execute_sinh_opascendclacl_op_compilernnopbasestdc
)install(TARGETS execute_sinh_op DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})
最后修改run.sh脚本中关于路径的部分。
修改完成后就可以执行run.sh脚本进行单算子API调用了。
INFO: acl executable run success!
input1: [ 0.468 -0.2585 -3.066 ... 0.9136 -1.117 -1.368 ]
output: [ 0.485 -0.2615 -10.71 ... 1.047 -1.365 -1.837 ]
-------------golden is :
golden: [ 0.4854 -0.2615 -10.71 ... 1.046 -1.364 -1.837 ]
INFO: compare golden data success!
出现上述提示证明算子通过验证。 5 Ascend C学习资源
Ascend C配套丰富的学习资料包括教程文档、交流社区、案例代码等这些资源将帮助您理解Ascend C编程语言的各种概念和技巧为您的自主学习提供便利。 Ascend C学习资源汇聚页Ascend C-昇腾社区 Ascend C官方文档昇腾社区-官网丨昇腾万里 让智能无所不及 Ascend C视频课程入门昇腾社区-官网丨昇腾万里 让智能无所不及 Ascend C视频课程进阶昇腾社区-官网丨昇腾万里 让智能无所不及 Ascend C论坛https://www.hiascend.com/forum/forum-0163125572293226003-1.html Ascend C sample开发者分享 | Ascend C算子开发及单算子调用
