【Ascend C算子开发(入门)】——Ascend C编程模式与范式

Ascend C编程模型与范式

1.并行计算架构抽象

Ascend C编程开发的算子是运行在AI Core上的,所以我们需要了解一下AI Core的结构。AI Core主要包括计算单元、存储单元、搬运单元。

  • 计算单元包括了三种计算资源:Scalar计算单元(执行标量计算);Cube计算单元(矩阵计算);Vector计算单元(向量运算)
  • 搬运单元主要负责在Global Memory 和Local Memory之间搬运数据
  • 包括内部存储(Local Memory)和外部存储(Global Memory)

数据在这些单元之间存储和计算,涉及到三种流:异步指令流、同步信号流、计算数据流。

  • 异步指令流:指计算单元和搬运单元之间异步执行接收到的指令序列。
  • 同步信号流:保证不同指令按照正确的逻辑关系执行
  • 计算数据流:指搬运单元把数据搬运到Local Memory,把处理好的数据搬运回Global Memory的过程。

AI Core的内部架构图如下:
image.png


2.SPMD编程模型介绍

SPMD(Single Program, Multiple Data)模型是一种并行编程模型,用于同时处理多个数据元素的相同程序。在Ascend C中,SPMD模型用于编写并行计算任务,以便充分利用Ascend AI处理器的并行计算能力。

SPMD模型的要点如下:

  • Single Program:SPMD模型意味着编写的程序是相同的,不会针对不同的数据元素而改变。这个程序会在不同的数据上执行,但代码本身是相同的。这有助于提高代码的可维护性和复用性。每个核上唯一区别是block_idx不同。

  • Multiple Data:程序会同时处理多个数据元素,这些数据元素通常存储在数组或张量中。每个数据元素都会被相同的程序逐一处理,从而实现并行性。


3.核函数编写及调用

Ascend C核函数是一种用于编写高性能并行计算任务的特定函数,是算子设备侧入口。

3.1核函数定义

8C18172DD5E2AA134AC39BB340E7592E.png

主要包括三个参数:函数类型限定符、核函数名、参数列表

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

2.指针入参变量需要增加变量类型限定符__gm__。表明该指针变量指向Global Memory上某处内存地址。

3.核函数使用内核调用符<<<…>>>这种语法形式,来规定核函数的执行配置:
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
解释每个参数的意思:

  • blockDim:规定了核函数将会在几个核上执行
  • l2ctrl:暂时设置为固定值nullptr,开发者无需关注
  • stream:类型为aclrtStream

昨天做了一个关于Ascend C加法算法的算子开发实验,代码地址:Gitee代码仓库
在这个Add文件中,算子开发的核心代码在 add_custom.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(); }
这段代码使用__global__ __aicore__函数类型限定符表明这个核函数将在AI Core上执行
void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z):这是核函数的声明,接受三个GM_ADDR参数,分别命名为x、y和z。GM_ADDR是指向通用内存(GM,General Memory)的指针类型,表明这个核函数将操作通用内存中的数据。

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); }

<<<blockDim, l2ctrl, stream>>>:这是CUDA执行配置,它指定了核函数 add_custom 的执行方式。blockDim 表示使用多少个CUDA线程块,l2ctrl 和 stream 表示与线程块配置和流相关的信息。这些参数通常用于控制CUDA核函数的执行方式和资源配置。

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/114147.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

外骨骼机器人和人形机器人概览

前言&#xff1a;一点思考 外骨骼机器人和人形机器人都曾随着一些爆品的出现火热过一段时间&#xff0c;但总感觉当前技术条件还不成熟&#xff0c;真正能落地的应用场景不多。马斯克在擎天柱发布会上被问到人形机器人的落地与前景问题时并没有给出明确答案&#xff0c;只是用…

c++ pcl点云变换骨架枝干添加树叶源码实例

程序示例精选 c pcl点云变换骨架枝干添加树叶源码实例 如需安装运行环境或远程调试&#xff0c;见文章底部个人QQ名片&#xff0c;由专业技术人员远程协助&#xff01; 前言 这篇博客针对《c pcl点云变换骨架枝干添加树叶源码实例》编写代码&#xff0c;代码整洁&#xff0c;…

2023年9月青少年软件编程(C 语言) 等级考试试卷(七级)

青少年软件编程&#xff08;C/C&#xff09;7级等级考试真题试卷&#xff08;2023年9月&#xff09; 编程题第 1 题 红与黑&#xff08;2023.9&#xff09; 有一间长方形的房子&#xff0c;地上铺了红色、黑色两种颜色的正方形瓷砖。你站在其中一块黑色的瓷砖上&#xff0c…

自然语言处理---Tr ansformer机制详解之Transformer结构

1 Encoder模块 1.1 Encoder模块的结构和作用 经典的Transformer结构中的Encoder模块包含6个Encoder Block.每个Encoder Block包含一个多头自注意力层&#xff0c;和一个前馈全连接层. 1.2 Encoder Block 在Transformer架构中&#xff0c;6个一模一样的Encoder …

Vue Router 刷新当前页面

Vue项目, 在实际工作中, 有些时候需要在 加载完某些数据之后对当前页面进行刷新, 以期 onMounted 等生命周期函数, 或者 数据重新加载. 总之是期望页面可以重新加载一次. 目前总结有三种途径可实现以上需求: 一, reload 直接刷新页面 window.location.reload(); $router.go(…

My Code Style

受到 _slb 大佬的启发&#xff0c;~~&#xff08;趁着集训休假开摆&#xff09;~~也打算写一篇关于自己码风的文章 QwQ。 自认为自己码风很好 头文件 万能头忠实拥护者&#xff0c;纵使它在变量名上虐我千百遍&#xff0c;我仍然要坚定地使用它&#xff01; 变量 数组一律…

【JavaScript】深入浅出理解事件循环

1. 浏览器的进程模型 1.1 进程 程序运行需要有它自己专属的内存空间&#xff0c;可以把这块内存空间简单的理解为进程。 每个应用至少有一个进程&#xff0c;进程之间相互独立&#xff0c;即使要通信&#xff0c;也需要双方同意。 1.2 线程 有了进程后&#xff0c;就可以运…

【广州华锐互动】VR营销心理学情景模拟培训系统介绍

在高度竞争的汽车市场中&#xff0c;销售人员需要具备强大的专业知识、引人入胜的销售技巧&#xff0c;以及敏锐的市场洞察力。然而&#xff0c;传统的培训方式往往无法满足这些需求&#xff0c;因为它们往往忽略了实践的重要性。 为了解决这个问题&#xff0c;许多公司开始采用…

老卫带你学---leetcode刷题(215. 数组中的第K个最大元素)

215. 数组中的第K个最大元素 问题&#xff1a; 给定整数数组 nums 和整数 k&#xff0c;请返回数组中第 k 个最大的元素。 请注意&#xff0c;你需要找的是数组排序后的第 k 个最大的元素&#xff0c;而不是第 k 个不同的元素。 你必须设计并实现时间复杂度为 O(n) 的算法解…

TCP/IP(十九)TCP 实战抓包分析(三)TCP 第一次握手 SYN 丢包

一 TCP 三次握手异常情况实战分析 说明&#xff1a; 本文是TCP 三次握手异常系列之一 ① 异常场景 接下里我用三个实验案例,带大家一起探究探究这三种异常关注&#xff1a; 如何刻意练习模拟上述场景 以及 wireshark现象 ② 实验环境 ③ 实验一&#xff1a;TCP 第一次握…

Python —— UI自动化之使用JavaScript进行元素点亮、修改、点击元素

1、JavaScript点亮元素 在控制台通过JavaScript语言中对元素点亮效果如下&#xff1a; 将这个语句和UI自动化结合&#xff0c;代码如下&#xff1a; locator (By.ID,"kw") # 是元组类型 web_element WebDriverWait(driver,5,0.5).until(EC.visibility_of_eleme…

Arduino驱动BMA220三轴加速度传感器(惯性测量传感器篇)

目录 1、传感器特性 2、硬件原理图 3、驱动程序 BMA220的三轴加速度计是一款具有I2C接口的超小型三轴低g加速度传感器断路器,面向低功耗消费市场应用。它可以测量3个垂直轴的加速度,从而在手机、手持设备、计算机外围设备、人机界面、虚拟现实功能和游戏控制器中感知倾斜、…

利用CSRF或XSS攻击网站的例子

利用 CSRF 攻击网站的简单示例&#xff1a; 假设有一个在线银行应用&#xff0c;用户可以在其中执行转账操作。用户登录后&#xff0c;系统会生成一个包含转账信息的表单&#xff0c;用户需要填写表单来发起转账。这个表单如下所示&#xff1a; <form action"https:/…

Python数据分析之numpy的使用

作者将自己学习numpy的代码笔记分享一下&#xff0c;想学哪一部分就把哪一部分的注释去掉即可&#xff0c;便于初学者学习和复习&#xff1a; import numpy as np主要是根据数据构建算法可能用到的矩阵&#xff0c;对于矩阵可以进行相应的处理变换#将一个列表输出成一维向量&a…

王道计算机考研 操作系统学习笔记 + 完整思维导图篇章五: IO管理

目录 IO设备的基本概念和分类 IO设备的分类 按使用特性分类 按传输速率分类 按信息交换单位分类 IO控制器 l/O设备的电子部件&#xff08;I/O控制器&#xff09; l/O控制器的组成 内存映像I/o vs.寄存器独立编址 IO控制方式 程序直接控制方式 中断驱动方式 DMA方式 ​编辑通…

java1.8流的新特性使用

案例描述 今天跟着黑马程序员的视频&#xff0c;完成“瑞吉外卖”项目的菜品信息管理模块的时候&#xff0c;遇到了一个比较陌生的写法 用到了Java8的新特性 stream().map((item) -> {}).collect() List<DishDto> collect records.stream().map((item) -> {DishDt…

Linux网络编程杂谈(聊聊网络编程背后的故事)

数据是如何传输到物理网络上的&#xff1f; 以TCP为例&#xff0c;当 TCP 决定发送数据时&#xff0c;这些数据需要经过多个处理阶段才能真正被传输到物理网络。其中一个关键步骤是将数据移动到网络接口卡 (NIC)。以下是这个过程的详细描述&#xff1a; 数据序列化: TCP 会为要…

10.17七段数码管单个多个(部分)

单个数码管的实现 第一种方式 一端并接称为位码&#xff1b;一端分别接收电平信号以控制灯的亮灭&#xff0c;称为段码 8421BCD码转七段数码管段码是将BCD码表示的十进制数转换成七段LED数码管的7个驱动段码&#xff0c; 段码就是LED灯的信号 a为1表示没用到a&#xff0c;a为…

文件读取结束的判定

大家好啊&#xff0c;我们今天来补充文件操作的读取结束的判定。 被错误使用的feof 牢记&#xff1a;在文件读取过程中&#xff0c;不能用feof函数的返回值直接用来判断文件的是否结束而是应用于当文件读取结束的时候&#xff0c;判断是读取失败结束&#xff0c;还是遇到文件尾…

快速排序(sort用法)

在头文件#include<algorithm>中 for (int i 0; i < m; i)cin >> arr[i];sort(arr, arr m);for (int i 0; i < m; i)cout << arr[i]; sort()函数可以对给定区间所有元素进行排序。 它有三个参数sort(begin, end, cmp)&#xff0c; 其中begin为指向…