首页 > 其他分享 >Ascend C算子开发指南2

Ascend C算子开发指南2

时间:2024-07-28 23:30:17浏览次数:13  
标签:指南 void Ascend ACL 算子 inputByteSize input CHECK

Ascend C算子开发指南
Ascend C的特点
C/C++原生编程:Ascend C原生支持C和C++标准规范。
屏蔽硬件差异:编程模型屏蔽了硬件差异,提高了代码的通用性。
API封装:类库API封装,既保证易用性,又兼顾高效性。
孪生调试:支持在CPU侧模拟NPU侧的行为,便于调试。
开发基本流程
环境准备:

安装CANN开发套件包,根据机器CPU架构下载对应的版本。
示例(AArch64架构):
bash
复制代码
wget -O Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run <下载链接>
chmod +x Ascend-cann-toolkit_8.0.RC1.alpha002_linux-x86_64.run
./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-x86_64.run --check
sudo ./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-x86_64.run --install
source /usr/local/Ascend/ascend-toolkit/set_env.sh
算子分析:

分析算子的数学表达式、输入输出数据类型和计算逻辑。
例如,Add算子的数学表达式为 $z = x + y$,输入输出数据类型为half(float16),支持的shape为(8, 2048)。
核函数开发(以Add算子为例):

获取样例代码目录quick-start,依次开发add_custom.cpp、main.cpp、gen_data.py三个文件。

核函数实现(add_custom.cpp):

cpp
复制代码
extern "C" global aicore void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
KernelAdd op;
op.Init(x, y, z);
op.Process();
}

void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z) {
add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}
算子类实现(KernelAdd):

cpp
复制代码
class KernelAdd {
public:
aicore inline KernelAdd() {}
aicore inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
// 初始化代码
}
aicore inline void Process() {
// 核心处理函数
}
private:
// 各阶段函数定义
aicore inline void CopyIn(int32_t progress) {}
aicore inline void Compute(int32_t progress) {}
aicore inline void CopyOut(int32_t progress) {}
private:
TPipe pipe;
TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
GlobalTensor xGm, yGm, zGm;
};
Process函数:

cpp
复制代码
aicore inline void Process() {
constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;
for (int32_t i = 0; i < loopCount; i++) {
CopyIn(i);
Compute(i);
CopyOut(i);
}
}
CopyIn函数:

cpp
复制代码
aicore inline void CopyIn(int32_t progress) {
LocalTensor xLocal = inQueueX.AllocTensor();
LocalTensor yLocal = inQueueY.AllocTensor();
DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
inQueueX.EnQue(xLocal);
inQueueY.EnQue(yLocal);
}
Compute函数:

cpp
复制代码
aicore inline void Compute(int32_t progress) {
LocalTensor xLocal = inQueueX.DeQue();
LocalTensor yLocal = inQueueY.DeQue();
LocalTensor zLocal = outQueueZ.AllocTensor();
Add(zLocal, xLocal, yLocal, TILE_LENGTH);
outQueueZ.EnQue(zLocal);
inQueueX.FreeTensor(xLocal);
inQueueY.FreeTensor(yLocal);
}
CopyOut函数:

cpp
复制代码
aicore inline void CopyOut(int32_t progress) {
LocalTensor zLocal = outQueueZ.DeQue();
DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
outQueueZ.FreeTensor(zLocal);
}
运行验证(main.cpp):

CPU侧验证:

cpp
复制代码
// 初始化内存并调用核函数
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);
AscendC::SetKernelMode(KernelMode::AIV_MODE);
ICPU_RUN_KF(add_custom, blockDim, x, y, z);
WriteFile("./output/output_z.bin", z, outputByteSize);
AscendC::GmFree((void *)x);
AscendC::GmFree((void *)y);
AscendC::GmFree((void *)z);
NPU侧验证:

cpp
复制代码
// 初始化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));
// 分配内存并进行数据拷贝
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));
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));
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(blockDim, nullptr, stream, xDevice, yDevice, zDevice);
CHECK_ACL(aclrtSynchronizeStream(stream));
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(aclrtDestroyContext(context));
CHECK_ACL(aclrtResetDevice(deviceId));
CHECK_ACL(aclFinalize());
数据生成(gen_data.py):

python
复制代码
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()
运行验证:

设置环境变量:
bash
复制代码
export ASCEND_HOME_DIR=/usr/local/Ascend/ascend-toolkit/latest
执行脚本:
bash
复制代码
bash run.sh <soc_version> <run_mode>
通过以上步骤,即可完成Ascend C算子的开发和验证。

标签:指南,void,Ascend,ACL,算子,inputByteSize,input,CHECK
From: https://www.cnblogs.com/taixian/p/18329143

相关文章

  • Ascend C算子开发指南
    AscendC的特点C/C++原生编程:AscendC原生支持C和C++标准规范。屏蔽硬件差异:编程模型屏蔽了硬件差异,提高了代码的通用性。API封装:类库API封装,既保证易用性,又兼顾高效性。孪生调试:支持在CPU侧模拟NPU侧的行为,便于调试。开发基本流程环境准备:安装CANN开发套件包,根据机器CPU架......
  • 昇思25天学习打卡营第16天|GAN 图像生成指南:数据集和模型训练手册
    目录MindSpore环境配置、MNIST数据集下载及处理展开。数据集可视化隐码构造模型构建模型训练效果展示模型推理MindSpore环境配置、MNIST数据集下载及处理展开。        首先,通过命令行操作安装特定版本的MindSpore库,并查看其版本。接着,从指定URL......
  • 提高 C# 的生产力:C# 13 更新完全指南
    提高C#的生产力:C#13更新完全指南 前言#预计在2024年11月,C#13将与.NET9一起正式发布。今年的C#更新主要集中在 refstruct 上进行了许多改进,并添加了许多有助于进一步提高生产力的便利功能。本文将介绍预计将在C#13中添加的功能。注意:目前C#13还未......
  • 利用python和工具变量法精确估计价格对销量的影响:解决内生性问题的实战指南
    目录1.引言2.工具变量法简介3.案例背景3.1背景信息3.2具体行动3.3数据收集4.实现思路5.示例代码5.1构建测试数据5.2进行工具变量法分析5.2.1一阶段回归5.2.2二阶段回归5.2.3验证工具变量5.3结论与应用5.3.1结论5.3.2实际应用6下一步思考6.1进一步验证和......
  • 鸣潮游戏错误126:加载x3daudio1_7.dll失败的全面解析与修复指南
    在畅玩鸣潮游戏时,不少玩家可能会遭遇错误代码「126」,提示“加载x3daudio1_7.dll失败,该文件缺失或损坏”。这个问题看似棘手,实则有迹可循,通过本文,我们将深入探讨其成因,并提供详细的解决步骤,帮助你重拾游戏乐趣。x3daudio1_7.dll是什么?x3daudio1_7.dll是一个与DirectX音频组件......
  • Postman中的代理艺术:配置与使用指南
    Postman中的代理艺术:配置与使用指南在API开发和测试过程中,代理服务器常用于捕获、检查、修改请求和响应。Postman作为一个流行的API开发工具,内置了代理服务器功能,使得测试人员可以方便地查看和修改通过代理的流量。本文将详细介绍如何在Postman中配置和使用代理服务器。代......
  • Django Web开发:构建强大RBAC权限管理系统的实战指南
    文章目录前言一、rbac基于角色的权限管理1.acl基于用户的权限管理2.rbac基于角色的权限管理二、应用示例1.配置角色资源a.分析表b.核心逻辑c.使用transfer在前端实现资源配置d.页面效果2.登录时获取对应权限a.员工登录b.中间件c.前端请求d.效果图3.前端-路由守卫......
  • pytorch中自定义onnx新算子并导出为onnx
    importtorchfromtorch.autogradimportFunctionimporttorch.onnx#Step1:DefinecustomPyTorchoperatorclassMyCustomOp(Function):@staticmethoddefforward(ctx,input):returninput+1@staticmethoddefsymbolic(g,input):......
  • 洛谷题单指南-前缀和差分与离散化-P1496 火烧赤壁
    原题链接:https://www.luogu.com.cn/problem/P1496题意解读:给定n个区间[a,b),计算所有区间覆盖的总长度。解题思路:方法1、离散化先思考一种比较直观的思路:既然要计算多个区间覆盖的总长度,可以枚举每一个区间[a,b),通过一个桶数组来标记区间中所有的点f[x]=1,最终统计所有为1的......
  • 如何选择合适的Bug跟踪软件?终极指南
    国内外主流的10款BUG管理软件对比:PingCode、Worktile、禅道(ZenTao)、Bugzilla、Tapd、CODING、Teambition、Testin、Tower、乐道。在软件开发的世界里,管理和跟踪Bug是一个让许多开发者头疼的问题。选择一个合适的Bug管理工具不仅能提升开发效率,还能大大减少因错误管理导致的延误......