课堂AScend C知识简记
AIcore三个计算单元Cube矩阵计算单元，Vector向量计算，Scalar标量计算
C/C++标准规范，自动化流水并行调度，结构化核函数编程，CPU/NPU孪生调试
三类api：计算类，同步类，搬运类
达芬奇架构主要组成部分：计算单元，存储单元，控制单元
三个流：异步指令流（cpu可以进行标量计算，cpu可以是scalar计算单元），同步指令流，计算数据流（DMA搬运单元实现G->L
L->G的数据传输)
算子可以看作一种函数
张量是存储算子输入数据与输出数据的容器
张量描述符：名称，形状，数据类型，数据排布格式
数据排布格式：
N:batch数量，例如图像的数目
H:高度
W:宽度
C:channel,特征图通道
device模块：负责指定计算运行的真实设备。
cpu->host内存
核函数跑在aicore上，所有核函数都返回void类型，返回需要通过参数传递
核函数可以被<<<>>>调用
blockDim规定了核函数将会在几个核上运行
l2ctrl，保留参数
stream：执行任务流
逻辑位置（初始化时使用）
for循环命令的指令输入放到队列中，三个队列指令流水执行，
核内切分，流水任务：stage1~n（流水优化）同一时间多个stage同时处理
double buffer队列两块数据，实现队列的流水
矢量逻辑位置
VECIN搬入数据存放位置
VECOUT搬出数据存放位置
VECCALC临时数据存放方式
BUFFER_NUM固定值为二
详细笔记
一、Ascend C and CANN
1.Ascend C是一种使用C/C++作为前端语言的算子开发工具，通过四层接口抽象、并行编程范式、孪生调试等技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。
2.CANN
CANN是华为针对AI场景推出的异构计算架构，对上支持多种AI框架，对下服务AI处理器与编程，发挥承上启下的关键作用，是提升昇腾AI处理器计算效率的关键平台。
二·、Ascend优点特性
Ascend C继承了TIK数据操作灵活的优点，除此之外，更多优势如下：
（1）C/C++原语编程
（2）编程模型屏蔽硬件差异，编程范式提高开发效率
（3）多层级API封装，从简单到灵活，兼顾易用与高效
（4）调试方法简单
三、核函数
1.基本定义
核函数（Kernel Function）是Ascend C算子设备侧的入口。Ascend C允许用户使用核函数这种C/C++函数的语法扩展来管理设备侧的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。
核函数是直接在设备侧执行的代码。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，当核函数被调用时，多个核将并行执行同一个计算任务。
2.编写规范
__global__ __aicore__ void kernel_name(argument list);
a.使用函数限定符
使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备侧AI Core上执行。
b.使用变量类型限定符
指针入参变量统一的类型定义为__gm__ uint8_t*。用户可统一使用uint8_t类型的指针，并在使用时转化为实际的指针类型；亦可直接传入实际的指针类型。
c.其他规则
必须具有void返回类型；使用extern “C”；仅支持入参为指针类型或C/C++内置数据类型(Primitive Data Types)，如：half* s0、float* s1、int32_t c。
3.调用方式
a.核函数使用内核调用符<<<...>>>这种语法形式，来规定核函数的执行配置。
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
blockDim，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑ID，表现为内置变量block_idx，编号从0开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用。
l2ctrl，保留参数，暂时设置为固定值nullptr。
stream，类型为aclrtStream，stream是一个任务队列，应用程序通过stream来管理任务的并行。
b.核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机侧。
强制主机侧程序等待所有核函数执行完毕的API（阻塞应用程序运行，直到指定Stream中的所有任务都完成，同步接口）为aclrtSynchronizeStream。
四、多级API的设计
基本概念
矢量计算指令接口，能够启动AI Core中的Vector单元执行计算。为了降低开发者的使用门槛，指令按照由易到难，分成了3级到0级接口。其中3级接口最为简单，0级接口最为复杂。
a. 3级接口，运算符重载，支持+, -, *, /, |, &, ^, >, < , >=, <=，!=，==实现2级接口的简化表达。
b. 2级接口，针对源操作数srcLocal的连续COUNT个数据进行计算，并连续写入目的操作数dstLocal，提供了一维Tensor的连续COUNT个数据的计算支持。
c. 0级接口，是最底层的开发接口，可以完整发挥硬件优势的计算API，可以进行非连续的计算该功能可以充分发挥CANN系列芯片的强大功能指令，支持对每个操作数的Block stride，Repeat stride，MASK的操作，允许用户使用诸多的通用参数来定制化所需要的操作。
五、中级认证的SinhCustom操作
1.将试题代码拷贝到开发板环境
2.修改搬入搬出函数，将搬入数据记为下x，搬出数据记为z，修改搬入搬出逻辑
将CarryIn修改如下
将CarryOut的修改代码如下
3.修改Compute逻辑，实现Sinh的具体计算，如下图
4.使用msopgen创建算子工程，分别在op_host和op_kernel分别进行处理
5.使用Aclnn方法进行调用
6.进行正确性验证
六、关于矩阵
Ascend C提供一组Matmul高阶API，方便用户快速实现Matmul矩阵乘法的运算操作。
MatMul的计算公式为：C = A * B + Bias，其示意图如下。
- A、B为源操作数，A为左矩阵，形状为[M, K]；B为右矩阵，形状为[K, N]。
- C为目的操作数，存放矩阵乘结果的矩阵，形状为[M, N]。
- Bias为矩阵乘偏置，形状为[1, N]。对A*B结果矩阵的每一行都采用该bias进行偏置。（深度学习时设计的偏置值）
图1 Matmul矩阵乘示意图
实现MatMul矩阵乘运算的具体步骤如下：
（1）创建Matmul对象。
（2）初始化操作
（3）设置左矩阵A、右矩阵B、Bias。
（4）完成矩阵乘操作。
（5）结束矩阵乘操作。
点赞
收藏
回复
分享
举报
课堂AScend C知识简记
AIcore三个计算单元Cube矩阵计算单元，Vector向量计算，Scalar标量计算
C/C++标准规范，自动化流水并行调度，结构化核函数编程，CPU/NPU孪生调试
三类api：计算类，同步类，搬运类
达芬奇架构主要组成部分：计算单元，存储单元，控制单元
三个流：异步指令流（cpu可以进行标量计算，cpu可以是scalar计算单元），同步指令流，计算数据流（DMA搬运单元实现G->L
L->G的数据传输)
算子可以看作一种函数
张量是存储算子输入数据与输出数据的容器
张量描述符：名称，形状，数据类型，数据排布格式
数据排布格式：
N:batch数量，例如图像的数目
H:高度
W:宽度
C:channel,特征图通道
device模块：负责指定计算运行的真实设备。
cpu->host内存
核函数跑在aicore上，所有核函数都返回void类型，返回需要通过参数传递
核函数可以被<<<>>>调用
blockDim规定了核函数将会在几个核上运行
l2ctrl，保留参数
stream：执行任务流
逻辑位置（初始化时使用）
for循环命令的指令输入放到队列中，三个队列指令流水执行，
核内切分，流水任务：stage1~n（流水优化）同一时间多个stage同时处理
double buffer队列两块数据，实现队列的流水
矢量逻辑位置
VECIN搬入数据存放位置
VECOUT搬出数据存放位置
VECCALC临时数据存放方式
BUFFER_NUM固定值为二
详细笔记
一、Ascend C and CANN
1.Ascend C是一种使用C/C++作为前端语言的算子开发工具，通过四层接口抽象、并行编程范式、孪生调试等技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。
2.CANN
CANN是华为针对AI场景推出的异构计算架构，对上支持多种AI框架，对下服务AI处理器与编程，发挥承上启下的关键作用，是提升昇腾AI处理器计算效率的关键平台。
二·、Ascend优点特性
Ascend C继承了TIK数据操作灵活的优点，除此之外，更多优势如下：
（1）C/C++原语编程
（2）编程模型屏蔽硬件差异，编程范式提高开发效率
（3）多层级API封装，从简单到灵活，兼顾易用与高效
（4）调试方法简单
三、核函数
1.基本定义
核函数（Kernel Function）是Ascend C算子设备侧的入口。Ascend C允许用户使用核函数这种C/C++函数的语法扩展来管理设备侧的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。
核函数是直接在设备侧执行的代码。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，当核函数被调用时，多个核将并行执行同一个计算任务。
2.编写规范
__global__ __aicore__ void kernel_name(argument list);
a.使用函数限定符
使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备侧AI Core上执行。
b.使用变量类型限定符
指针入参变量统一的类型定义为__gm__ uint8_t*。用户可统一使用uint8_t类型的指针，并在使用时转化为实际的指针类型；亦可直接传入实际的指针类型。
c.其他规则
必须具有void返回类型；使用extern “C”；仅支持入参为指针类型或C/C++内置数据类型(Primitive Data Types)，如：half* s0、float* s1、int32_t c。
3.调用方式
a.核函数使用内核调用符<<<...>>>这种语法形式，来规定核函数的执行配置。
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
blockDim，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑ID，表现为内置变量block_idx，编号从0开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用。
l2ctrl，保留参数，暂时设置为固定值nullptr。
stream，类型为aclrtStream，stream是一个任务队列，应用程序通过stream来管理任务的并行。
b.核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机侧。
强制主机侧程序等待所有核函数执行完毕的API（阻塞应用程序运行，直到指定Stream中的所有任务都完成，同步接口）为aclrtSynchronizeStream。
四、多级API的设计
基本概念
矢量计算指令接口，能够启动AI Core中的Vector单元执行计算。为了降低开发者的使用门槛，指令按照由易到难，分成了3级到0级接口。其中3级接口最为简单，0级接口最为复杂。
a. 3级接口，运算符重载，支持+, -, *, /, |, &, ^, >, < , >=, <=，!=，==实现2级接口的简化表达。
b. 2级接口，针对源操作数srcLocal的连续COUNT个数据进行计算，并连续写入目的操作数dstLocal，提供了一维Tensor的连续COUNT个数据的计算支持。
c. 0级接口，是最底层的开发接口，可以完整发挥硬件优势的计算API，可以进行非连续的计算该功能可以充分发挥CANN系列芯片的强大功能指令，支持对每个操作数的Block stride，Repeat stride，MASK的操作，允许用户使用诸多的通用参数来定制化所需要的操作。
五、中级认证的SinhCustom操作
1.将试题代码拷贝到开发板环境
2.修改搬入搬出函数，将搬入数据记为下x，搬出数据记为z，修改搬入搬出逻辑
将CarryIn修改如下
将CarryOut的修改代码如下
3.修改Compute逻辑，实现Sinh的具体计算，如下图
4.使用msopgen创建算子工程，分别在op_host和op_kernel分别进行处理
5.使用Aclnn方法进行调用
6.进行正确性验证
六、关于矩阵
Ascend C提供一组Matmul高阶API，方便用户快速实现Matmul矩阵乘法的运算操作。
MatMul的计算公式为：C = A * B + Bias，其示意图如下。
图1 Matmul矩阵乘示意图
实现MatMul矩阵乘运算的具体步骤如下：
（1）创建Matmul对象。
（2）初始化操作
（3）设置左矩阵A、右矩阵B、Bias。
（4）完成矩阵乘操作。
（5）结束矩阵乘操作。