OrangePi AIpro 香橙派 昇腾 Ascend C算子开发 - HelloWorld

2024-08-31 06:52
文章标签 开发 算子 香橙 helloworld ascend orangepi aipro

本文主要是介绍OrangePi AIpro 香橙派 昇腾 Ascend C算子开发 - HelloWorld，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！

OrangePi AIpro 香橙派 昇腾 Ascend C算子开发 - HelloWorld

flyfish

Ascend C算子编程是SPMD（Single-Program Multiple-Data）编程
假设，从输入数据到输出数据需要经过3个阶段任务的处理（T1、T2、T3）。如下图所示，SPMD会启动一组进程，并行处理待处理的数据。对待处理数据切分，把切分后数据分片分发给不同进程处理，每个进程对自己的数据分片进行3个任务的处理。
在这里插入图片描述
具体到Ascend C编程模型中的应用，是将需要处理的数据被拆分并同时在多个计算核心（类比于上文介绍中的多个进程）上运行，从而获取更高的性能。多个AI Core共享相同的指令代码，每个核上的运行实例唯一的区别是block_idx不同，每个核通过不同的block_idx来识别自己的身份。block的概念类似于上文中进程的概念，block_idx就是标识进程唯一性的进程ID。并行计算过程的示意图如下图所示。
在这里插入图片描述
核函数（Kernel Function）是Ascend C算子设备侧实现的入口。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，当核函数被调用时，多个核都执行相同的核函数代码，具有相同的参数，并行执行。

Ascend C允许用户使用核函数这种C/C++函数的语法扩展来管理设备端的运行代码，用户在核函数中进行算子类对象的创建和其成员函数的调用，由此实现该算子的所有功能。核函数是主机端和设备端连接的桥梁。

一个Hello World例子展示Ascend C核函数（设备侧实现的入口函数）的基本写法和如何被调用的流程。

hello_world.cpp

#include "kernel_operator.h"extern "C" __global__ __aicore__ void hello_world()
{AscendC::printf("Hello World!!!\n");
}void hello_world_do(uint32_t blockDim, void *stream)
{hello_world<<<blockDim, nullptr, stream>>>();
}

main.cpp

#include "acl/acl.h"
extern void hello_world_do(uint32_t coreDim, void *stream);int32_t main(int argc, char const *argv[])
{aclInit(nullptr);int32_t deviceId = 0;aclrtSetDevice(deviceId);aclrtStream stream = nullptr;aclrtCreateStream(&stream);constexpr uint32_t blockDim = 8;hello_world_do(blockDim, stream);aclrtSynchronizeStream(stream);aclrtDestroyStream(stream);aclrtResetDevice(deviceId);aclFinalize();return 0;
}

HelloWorldSample例子

下载地址
进入源码目录执行

 bash run.sh -v Ascend310B4

结果

opType=hello_world, DumpHead: AIV-0, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!
opType=hello_world, DumpHead: AIV-1, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!
opType=hello_world, DumpHead: AIV-2, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!
opType=hello_world, DumpHead: AIV-3, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!
opType=hello_world, DumpHead: AIV-4, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!
opType=hello_world, DumpHead: AIV-5, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!
opType=hello_world, DumpHead: AIV-6, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!
opType=hello_world, DumpHead: AIV-7, CoreType=, block dim=8, total_block_num=8, block_remain_len=1048424, block_initial_space=1048576, rsv=0, magic=5aa5bccd
CANN Version: 901005402, TimeStamp: 20240821
Hello World!!!

在这里插入图片描述

extern "C" __global__ __aicore__ void hello_world()

核函数时需要遵循以下规则

使用函数类型限定符

除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__。

使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；
使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行：

__global__ __aicore__ void kernel_name(argument list);

编程中使用到的函数可以分为三类：核函数（device侧执行）、host侧执行函数、device侧执行函数（除核函数之外的）。三者的调用关系如下图所示：

host侧执行函数可以调用同类的host执行函数，也就是通用C/C++编程中的函数调用；也可以通过<<<>>>调用核函数。
device侧执行函数（除核函数之外的）可以调用同类的device侧执行函数。
核函数可以调用device侧执行函数（除核函数之外的）。
核函数（device侧执行）、host侧执行函数、device侧执行函数（除核函数之外的）调用关系图
在这里插入图片描述

使用变量类型限定符

指针入参变量需要增加变量类型限定符__gm__。表明该指针变量指向Global Memory上某处内存地址。
其他规则或建议

规则：核函数必须具有void返回类型。
规则：仅支持入参为指针或C/C++内置数据类型（Primitive data types），如：half* s0、float* s1、int32_t c。
建议：为了统一表达，建议使用GM_ADDR宏来修饰入参，GM_ADDR宏定义如下：

#define GM_ADDR __gm__ uint8_t*

使用GM_ADDR修饰入参的样例如下：

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)

这里统一使用uint8_t类型的指针，在后续的使用中需要将其转化为实际的指针类型。

代码解释 extern "C"

extern "C" 告诉编译器，不要对这些函数名进行C++的名称修饰，这样C语言的代码就可以正确地调用这些函数。
如果有一部分代码是用C编写的，另一部分代码是用C++编写的，extern "C" 会确保 my_c_function 以C语言的方式进行链接,C++编写的代码能够调用C语言编写的函数。
名称修饰是编译器生成唯一符号名的一种机制，目的是支持C++的高级功能，如函数重载。通过名称修饰，编译器确保每个函数或变量在链接阶段具有唯一性，避免命名冲突。

举个例子

考虑以下两个函数：

int add(int a, int b);
double add(double a, double b);

在C语言中，由于函数名称不能重载，这两个函数将会引起命名冲突。但是在C++中，编译器会将这两个函数分别转换为不同的符号名，例如（符号名称是编译器生成的，具体表示可能会不同）：

_add_int_int
_add_double_double

这些修饰后的名称在编译后的二进制文件中存储，使得它们可以在链接时区分开来。
使用 extern "C" 避免名称修饰如果希望C++函数能够被C代码调用，或者希望C++代码调用C语言的函数，需要使用extern "C"来告诉编译器不要对这些函数进行名称修饰。例如：

extern "C" void myFunction(int a);

在这种情况下，myFunction 的名字在编译后的二进制文件中将保持为 myFunction，而不会被修饰。

代码解释 extern void hello_world_do(uint32_t coreDim, void *stream);中的extern

在C++中，extern 关键字用于声明一个变量或函数是由其他文件定义的，而不是在当前文件中定义的。它告诉编译器这个函数的定义在另一个编译单元（例如另一个源文件）中，而不是在当前文件中。在提供的 main.cpp 文件中，extern void hello_world_do(uint32_t coreDim, void *stream); 这一行的作用是声明 hello_world_do 函数的存在，使得 main.cpp 文件可以调用这个函数，而不需要在 main.cpp 中定义它。

原理：

  1. 函数定义在另一个文件中 ：
    hello_world_do 函数实际上是在 hello_world.cpp 文件中定义的。为了在 main.cpp 中使用这个函数，编译器需要知道这个函数的签名（返回类型、参数类型等）。通过使用 extern，告诉编译器“这个函数在别的地方定义了，只需要知道它的签名就可以了”。

  2. 链接阶段的作用 ：
    编译器在编译 main.cpp 时，不需要知道 hello_world_do 函数的具体实现，只需要知道它的签名。而在链接阶段，链接器会把 hello_world.cpp 中的 hello_world_do 函数实现与 main.cpp 中的调用关联起来。

执行配置由3个参数决定：

blockDim，规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx来获取block_idx；
l2ctrl，保留参数，暂时设置为固定值nullptr，开发者无需关注；
stream，类型为aclrtStream，stream用于维护一些异步操作的执行顺序，确保按照应用程序中的代码调用顺序在device上执行。

在这里插入图片描述计算单元包括了三种基础计算资源：Cube计算单元、Vector计算单元和Scalar计算单元。
存储单元包括内部存储和外部存储：

AI Core的内部存储，统称为Local Memory，对应的数据类型为LocalTensor。由于不同芯片间硬件资源不固定，可以为UB、L1、L0A、L0B等。
AI Core能够访问的外部存储称之为Global Memory，对应的数据类型为GlobalTensor。

DMA（Direct Memory Access）搬运单元：负责在Global Memory和Local Memory之间搬运数据

这篇关于OrangePi AIpro 香橙派 昇腾 Ascend C算子开发 - HelloWorld的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！


http://www.chinasem.cn/article/1123199。

相关文章

基于Qt开发一个简单的OFD阅读器

基于Qt开发一个简单的OFD阅读器

《基于Qt开发一个简单的OFD阅读器》这篇文章主要为大家详细介绍了如何使用Qt框架开发一个功能强大且性能优异的OFD阅读器,文中的示例代码讲解详细,有需要的小伙伴可以参考一下... 目录摘要引言一、OFD文件格式解析二、文档结构解析三、页面渲染四、用户交互五、性能优化六、示例代码七、未来发展方向八、结论摘要
阅读更多...
在 VSCode 中配置 C++ 开发环境的详细教程

在 VSCode 中配置 C++ 开发环境的详细教程

《在VSCode中配置C++开发环境的详细教程》本文详细介绍了如何在VisualStudioCode（VSCode）中配置C++开发环境,包括安装必要的工具、配置编译器、设置调试环境等步骤,通... 目录如何在 VSCode 中配置 C++ 开发环境：详细教程1. 什么是 VSCode？2. 安装 VSCo
阅读更多...
C#图表开发之Chart详解

C#图表开发之Chart详解

《C#图表开发之Chart详解》C#中的Chart控件用于开发图表功能,具有Series和ChartArea两个重要属性,Series属性是SeriesCollection类型,包含多个Series对... 目录OverviChina编程ewSeries类总结OverviewC#中，开发图表功能的控件是Char
阅读更多...
鸿蒙开发搭建flutter适配的开发环境

鸿蒙开发搭建flutter适配的开发环境

《鸿蒙开发搭建flutter适配的开发环境》文章详细介绍了在Windows系统上如何创建和运行鸿蒙Flutter项目,包括使用flutterdoctor检测环境、创建项目、编译HAP包以及在真机上运... 目录环境搭建创建运行项目打包项目总结环境搭建1.安装 DevEco Studio NEXT IDE
阅读更多...
Python开发围棋游戏的实例代码（实现全部功能）

Python开发围棋游戏的实例代码（实现全部功能）

《Python开发围棋游戏的实例代码（实现全部功能）》围棋是一种古老而复杂的策略棋类游戏,起源于中国,已有超过2500年的历史,本文介绍了如何用Python开发一个简单的围棋游戏,实例代码涵盖了游戏的... 目录1. 围棋游戏概述1.1 游戏规则1.2 游戏设计思路2. 环境准备3. 创建棋盘3.1 棋盘类
阅读更多...
这15个Vue指令，让你的项目开发爽到爆

这15个Vue指令，让你的项目开发爽到爆

1. V-Hotkey 仓库地址: github.com/Dafrok/v-ho… Demo: 戳这里 https://dafrok.github.io/v-hotkey 安装: npm install --save v-hotkey 这个指令可以给组件绑定一个或多个快捷键。你想要通过按下 Escape 键后隐藏某个组件，按住 Control 和回车键再显示它吗？小菜一碟： <template
阅读更多...
Hadoop企业开发案例调优场景

Hadoop企业开发案例调优场景

需求 （1）需求：从1G数据中，统计每个单词出现次数。服务器3台，每台配置4G内存，4核CPU，4线程。 （2）需求分析： 1G / 128m = 8个MapTask；1个ReduceTask；1个mrAppMaster 平均每个节点运行10个 / 3台 ≈ 3个任务（4    3    3） HDFS参数调优 （1）修改：hadoop-env.sh export HDFS_NAMENOD
阅读更多...
嵌入式QT开发：构建高效智能的嵌入式系统

嵌入式QT开发：构建高效智能的嵌入式系统

摘要： 本文深入探讨了嵌入式 QT 相关的各个方面。从 QT 框架的基础架构和核心概念出发，详细阐述了其在嵌入式环境中的优势与特点。文中分析了嵌入式 QT 的开发环境搭建过程，包括交叉编译工具链的配置等关键步骤。进一步探讨了嵌入式 QT 的界面设计与开发，涵盖了从基本控件的使用到复杂界面布局的构建。同时也深入研究了信号与槽机制在嵌入式系统中的应用，以及嵌入式 QT 与硬件设备的交互，包括输入输出设
阅读更多...
OpenHarmony鸿蒙开发（ Beta5.0）无感配网详解

OpenHarmony鸿蒙开发（ Beta5.0）无感配网详解

1、简介 无感配网是指在设备联网过程中无需输入热点相关账号信息，即可快速实现设备配网，是一种兼顾高效性、可靠性和安全性的配网方式。 2、配网原理 2.1 通信原理 手机和智能设备之间的信息传递，利用特有的NAN协议实现。利用手机和智能设备之间的WiFi 感知订阅、发布能力，实现了数字管家应用和设备之间的发现。在完成设备间的认证和响应后，即可发送相关配网数据。同时还支持与常规Sof
阅读更多...
活用c4d官方开发文档查询代码

活用c4d官方开发文档查询代码

当你问AI助手比如豆包，如何用python禁止掉xpresso标签时候，它会提示到 这时候要用到两个东西。https://developers.maxon.net/论坛搜索和开发文档 比如这里我就在官方找到正确的id描述 然后我就把参数标签换过来
阅读更多...

Copyright China编程 23002807@qq.com

沪ICP备2023033072号-2