【2023 · CANN训练营第一季】TIK C++算子开发入门 第一章——TIK C++算子开发入门

news/2024/11/30 7:58:47/

1.TIK C++介绍

TIK C++是一种使用C/C++作为前端语言的算子开发工具,通过四层接口抽象、并行编程范式、孪生调试等技术,极大提高算子开发效率,助力AI开发者低成本完成算子开发和模型调优部署

使用TIK C++开发自定义算子的优势:
(1)C/C++原语编程
(2)编程模型屏蔽硬件差异,编程范式提高开发效率
(3)多层级API封装,从简单到灵活,兼顾易用与高效
(4)孪生调试,CPU侧模拟NPU侧的行为,可先在CPU侧调试

2.核函数

核函数(Kernel Function)是TIK C++算子设备侧的入口。TIK C++允许用户使用核函数这种C/C++函数的语法扩展来管理设备侧的运行代码,用户在核函数中实现算子逻辑的编写,例如自定义算子类及其成员函数以实现该算子的所有功能。核函数是主机侧和设备侧连接的桥梁
image.png
核函数是直接在设备侧执行的代码。在核函数中,需要为在一个核上执行的代码规定要进行的数据访问和计算操作,当核函数被调用时,多个核将并行执行同一个计算任务

2.1编写核函数

(1)使用函数类型限定符
除了需要按照C/C++函数声明的方式定义核函数之外,还要为核函数加上额外的函数类型限定符,包含__global__和__aicore__
使用__global__函数类型限定符来标识它是一个核函数,可以被<<<…>>>调用;使用__aicore__函数类型限定符来标识该核函数在设备侧AI Core上执行:
__global__ __aicore__ void kernel_name(argument list);
(2)使用变量类型限定符
为了方便:指针入参变量统一的类型定义为__gm__ uint8_t*
用户可统一使用uint8_t类型的指针,并在使用时转化为实际的指针类型;亦可直接传入实际的指针类型

其他规则
必须具有void返回类型
使用extern “C”
仅支持入参为指针类型或C/C++内置数据类型(Primitive Data Types),如:half* s0、float* s1、int32_t c

核函数的调用语句是C/C++函数调用语句的一种扩展
常见的C/C++函数调用方式是如下的形式:
function_name(argument list);
核函数使用内核调用符<<<…>>>这种语法形式,来规定核函数的执行配置:
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
blockDim,规定了核函数将会在几个核上执行,每个执行该核函数的核会被分配一个逻辑ID,表现为内置变量block_idx,编号从0开始,可为不同的逻辑核定义不同的行为,可以在算子实现中使用
l2ctrl,保留参数,暂时设置为固定值nullptr
stream,类型为aclrtStream,stream是一个任务队列,应用程序通过stream来管理任务的并行

3.接口介绍

常用数据定义:GlobalTensor
GlobalTensor用来存放Global Memory(外部存储)的全局数据
定义原型
template <typename T> class GlobalTensor { // 传入全局数据的指针,并手动设置一个buffer size,初始化GlobalTensor void SetGlobalBuffer(__gm__ T* buffer, uint32_t bufferSize);

LocalTensor
LocalTensor用于存放核上Local Memory(内部存储)的数据
定义原型
template <typename T> class LocalTensor { T GetValue(const uint32_t offset) const; // 获取 LocalTensor 中的某个值,返回 T 类型的立即数。 template <typename T1> void SetValue(const uint32_t offset, const T1 value) const; // 设置 LocalTensor 中的某个值。offset单位为element // 获取距原LocalTensor起始地址偏移量为offset的新LocalTensor,注意offset不能超过原有LocalTensor的size大小。offset单位为element LocalTensor operator[](const uint32_t offset) const; uint32_t GetSize() const; // 获取当前LocalTensor size大小

矢量计算指令接口
image.png


http://www.ppmy.cn/news/135387.html

相关文章

RNN基础概念

一、潜变量回归模型 使用潜变量 h t h_{t} ht​总结过去的信息 二、RNN 更新隐藏状态&#xff1a; h t φ ( W h h h t − 1 W h x x t − 1 b h ) h_{t}φ(W_{hh}h_{t-1}W_{hx}x_{t-1}b_{h}) ht​φ(Whh​ht−1​Whx​xt−1​bh​) 更新输出&#xff1a; o t W o h h t b…

创业老板野心有多大?我觉得可以搞个爱奇艺离线版试试……

黑客技术 点击右侧关注&#xff0c;了解黑客的世界&#xff01; Java开发进阶 点击右侧关注&#xff0c;掌握进阶之路&#xff01; Python开发 点击右侧关注&#xff0c;探讨技术话题&#xff01; 互联网行业人才辈出&#xff0c;特别是创业老板&#xff0c;更是有着无限的激情…

TCL爱奇艺电视TV+怎么安装第三方软件

TCL爱奇艺电视TV怎么安装第三方软件   很多用户买了TCL智能电视&#xff0c;首先要问到的是怎么安装直播软件&#xff0c;因为受政策影响&#xff0c;TCL智能电视本身是没有直播电视的。 准备一个sd卡&#xff0c;用电脑在sd卡上建一个tvmanager文件目录&#xff0c;再在这…

IMAX Enhanced:让沉浸式家庭影音娱乐体验不再抽象

IMAX所提供的观影体验瞬间可以打开观影者的视听神经&#xff0c;在一定程度上已经成为“视听盛宴”的一个具象化落地。对影音爱好者来说&#xff0c;自然也很希冀在电影的院线档期结束后仍有一个渠道能够重温这种观影体验&#xff0c;在家中也能享受到极具IMAX特色的观影畅爽感…

快录屏 v1.3.8

类型&#xff1a;影音播放 版本&#xff1a;v1.3.8 大小&#xff1a;13.7M 更新&#xff1a;2019/3/4 语言&#xff1a;简体 等级&#xff1a; 平台&#xff1a;安卓, 4.0以上 下载地址&#xff1a; 快录屏 v1.3.8&#xff08;1&#xff09; 快录屏 v1.3.8&#xff08;2&#x…

全网最详细的纪录片观看&下载指南

虽然当下短视频成了主流&#xff0c;愿意看纪录片这种冗长节目的人估计是少数了 但是我坚信纪录片的价值是巨大的&#xff0c;无论是帮助你开阔视野、认识世界、提升思维的维度… 尽管在信息洪流中算是比较小众的内容&#xff0c;阿虚还是想尽可能的传播这些有价值的东西 前…

三类主流影音播放器对比

三类主流影音播放器对比 本次调研主要针对的产品有 腾讯视频、暴风影音、迅雷影音 简单介绍 暴风影音 暴风影音是北京暴风科技有限公司推出的一款视频播放软件&#xff0c;该播放器兼容大多数的视频和音频格式。暴风影音播放的文件清晰&#xff0c;当有文件不可播时&#xff0c…

暴风陨落,再无影音​

暴风集团&#xff0c;迎来了它最终的归宿。 深交所在8月28日夜间发布公告称&#xff0c;决定暴风集团股票终止上市。 与刚刚退市的乐视一样&#xff0c;暴风集团也曾在主业不济、不足以撑起人的野心的时候&#xff0c;涉足了互联网电视、VR、体育等多个领域&#xff0c;还把触角…