embedded electronic systems for AI/ML 笔记第五篇

最新推荐文章于 2024-09-15 15:29:29 发布

无规则@lucio

最新推荐文章于 2024-09-15 15:29:29 发布

阅读量330

点赞数

分类专栏： AI/ML 复习笔记文章标签：人工智能 fpga开发

本文链接：https://blog.csdn.net/lsz92mjr/article/details/128259493

版权

AI/ML 复习笔记专栏收录该内容

7 篇文章 0 订阅

订阅专栏

一些嵌入式系统的人工智能和机器学习的平台介绍

主要平台
- 平行的可替代形式
- - ILP和TLP的弊端（drawbacks）
矢量处理器
GPU
GPU和矢量处理器的总结
GPU执行模型
FPGA
- Zynq UltraScale+ reconfigurable SOC
SOCs with ML accelerators
ASIC
- Google Tensor Processing Unit
- Google Edge TPU
总结

主要平台

这些平台主要有SISD（单指令单数据），SIMD（单指令多数据），MIMD（多指令多数据），MISD（多指令单数据）四种结构。
对于vector processor矢量处理器：主要有Intel Xeon Phi（SIMD）和Movidius等平台
对于GPU 主要有Nvidia Tesla（SIMD）和Xavier等平台
对于FPGA主要有Xilinx Virtex 和 Zynq等平台
对于SOC with ML acceleraors （又叫带机器学习加速器的片上系统）主要有Apple A14和HiSilicon Kirin等平台
对于ASIC主要有Google TPU和Edge TPU等平台。

平行的可替代形式

ILP（指令级）：super-scalar CPU or VLIM
TLP（线程级）：multi-threaded CPU
DP（矢量数据）：vector processor

ILP和TLP的弊端（drawbacks）

ILP的弊端：

超标量（super-scalar）：有太多的能量消耗和占用太多的面积，没有编译器支持情况下只能发现有限的并行。
过长指令字（Very-Long Instruction Word）：在每周期获取N指令和解码过程要消耗大量能量，这些指令可能包含大量NOP，密度差。

TLP的弊端：
在书写效率和纠正多线程代码上需要程序员付出更多的的工作量。

矢量处理器

结构：SIMD
缺点：通常可矢量化的代码很难被发现。
在这里插入图片描述
操作工作在N个数据的矢量上，数据存储在矢量寄存器中，使用矢量化pipelined ALU，使用矢量化LSU

Intel Xeon Phi

多核矢量处理器
参数：

61个X86 Atom 内核（每个内核有1.0-1.6GHz的时钟，较慢）：五个阶段顺序执行的流水线和每个核心有四个线程。
32个512-bit的矢量寄存器
512位（32字节）的矢量数据路径（vector datapath）：16个32位整数或单精度FP或者8个64位长整数或双精度浮点数。
具有分散/聚集（scatter/gather）和预取（prefetch）能力的LSU
1GHz6016*2 = 2TFLOP/s （1 MAC = 2 FLOP）
200-300 W ：10 GFLOP/J

结构：双向环形互连（Bidirectional ring interconnect），该结构也用于主存储器控制器中。

存储器结构

与host processor分离的存储器空间
12个或者16个GDDR5 存储器通道（5.5GT/s）

DDR DRAM

普及一下DDR是双倍数据速率（Double data rate），它的命令和地址在时钟上升沿激活，在两个时钟沿进行突发（burst）数据传输，所以增加了数据传输的性能而并没有要求让时钟比数据更快。
在这里插入图片描述

Movidius

Movidius Myriad 2 coprocessor

参数：和Xeon Phi不同的是嵌入了28nm的TSMC技术。
2个LEON（sparc v8）处理器（600 MHz）
12个SHAVE VLIW 矢量处理器（180MHz）
每个片上（on-chip）SRAM为128KB（一共1.5MB）
I/O 接口有SPI，I2C，I2S，SD，Ethernet和USB等
显示器和摄像头：MIPI，CIF，LCD
SIPP 图像信号处理加速器
结构
在这里插入图片描述

其中，每个SHAVE单元是一个带有SIMD的VLIM（视觉识别系统）
在200MHz频率下，每周期可执行多达5*12个FP指令，一共12GFLOP/s。

其中，

IAU（ integer Arithmetic Units），SAU（ Scalar Arithmetic Units），VAU（ Vector Arithmetic Units）每个周期可执行一个指令，SAU+VAU可多达1+4个32位FP指令，IAU+SAU+VAU可多大2+16个8位int指令。
CMU（Compare-and-Move Unit）可被用于产生predocates去使能VAU,IAU和SAU无分支指令的条件执行，也将数据在寄存器间移动。
BRU（Branch-and-Repeat Unit）可被用于实现固定迭代数循环（fixed iteration count loops）的零开销执行（zero-overhead execution）
LS（Load-Store Unit）实现在一周期内64位读和写操作。

一个SHAVE 2 核心的总体存储器带宽为（在频率为180MHz时）：
12.2 GB/s 的整数寄存器文件（有17个12位的端口）
8.6 GB/s的超标量寄存器文件（有12个4位端口）
34.6 gb/s的矢量寄存器文件（有12位16位端口）
1.9GB/s的片上SRAM（有两个8位端口）
1.4GB/s的片外（off-chip）封装（in-package）DRAM（有两个4位端口）

Movidius Myriad X coprocessor

参数：和上述相似的主机处理器接口和加速器（现在是700 MHz）
嵌入了16nm的TSMC技术
16个SHAVE VLIW 矢量处理器（200MHz）
每个片上SRAM有128kb（一共2MB）
神经计算引擎（Neural Compute Engine）能够达到1TIOP/s（8 bit int）
总峰值性能可达：4TIOP/s（8-bit int），4TIOP/J at 1W

总结

矢量处理器提供非常高的性能，每个操作的能量很低。因为简化了标量和视觉识别系统（VLIM）结构，由于SIMD所以共享了fetch/decode的开销（overhead）
对于DNN训练和迭代是一个很好的结构。
数据路径能够重新配置变量精度
寄存器文件和片上SRAM具有较大的带宽

GPU

全称为图形处理单元（Graphics Processing Unit），是一个多核处理器。
特点：每片有100-1000个核心，每个多核处理器比Xeon多核处理器更小更简单。每个处理器之间共享存储器，缓存空间和控制；大量使用SIMD，大量的多线程去隐藏存储器延时。

多核GPU

运行并行代码，主要是FP吞吐量。多个小核，通过多线程和手动管理片上系统的SRAM和寄存器来隐藏存储器延时。
在这里插入图片描述

CPU和GPU之间的较量

两者结构互补（complementary）
CPU对于大多数连续的，面向延时带有少量线程的代码更好。
GPU对于并行的，面向吞吐量代码更好，这些代码有数千个线程，需要更高的存储器带宽。

NVIDIA GT200 GPU

在这里插入图片描述
其中，一个FP32中有1024个寄存器，可多达2048个线程。

Nvidia GV100的多片HBM封装

高存储器带宽（high memory bandwidth）
被用于CPU，GPU，FPGA等
在这里插入图片描述

Nvidia Xavier

top layer

顶层的结构组成
在这里插入图片描述

OCTA-CORE Carmel

8核ARM8 CPU集群（cluster），每个小组有2MB的二级缓存。
在这里插入图片描述

Volta

英伟达Volta GPU，有8个SM，每个可以执行16个FP16 MAC/cycle，也能执行128个 int8 MAC/cycle
能执行：

22.6 TIOP/s int8
2.8 TFLOP/s FP16
1.4 TFLOP/s FP32

PVA

可编程视觉加速器（programmable vision accelerators）
使用通用机器视觉滤波器：Harris ，corner，FFT
包含：
ARM Cortex R5 core
两个 256位的7-issue（包含 2 scalar 2 vector 3 memory） VLIM 矢量处理器
328 bit 或者832bit 整数
1.7 TIOP/s int8
在这里插入图片描述

MM/DLA

深度学习加速器（deep learning accelerator）
11.7 TIOP/s int8
5.7 TFLOP/S fp16
立体和光学流发动机（Stereo and optical flow
engines）
12.8 TIOP/s int8
6.2 TIOP/s int16
存储器控制器（memory controller）
8个32位通道
127GB/s

在这里插入图片描述

汇总

22.6+1.7+11.7+12.8 ≈ 50 TIOP/s int8 (inference)
2.8+5.7 ≈ 8.5 TFLOP/s FP16 (inference)
1.4 TFLOP/s FP32 (training)
30 W
1600 TIOP/J int8
280 GFLOP/J FP16
46 GFLOP/J FP32
127 GB/s DDR4 bandwidth
2x2cm chip

GPU和矢量处理器的总结

在这里插入图片描述

GPU执行模型

CUDA和OpenCL
属于类C语言，但是补充了一下三点：

一些函数被标记为在加速器设备上运行
在Work Items（threads）和Work Groups（thread group）上隐式“doall”循环：这里doall的意思是所有的循环迭代都由程序员保证以任何顺序执行，没有干扰或者竞争。
显示地分配数组（和通常变量）到存储器层次结构（DRAM，SRAM，registers）。

关键原理（key principle）

用函数（内核）代替循环，在问题域地每一点上并行运行。

// execute n times in parallel
__kernel void
mul(__global const float *a,
__global const float *b,
__global float *c)
{
int i = get_global_id(0);
c[i] = a[i] * b[i];
}

而不是再用C语言地for循环。

编程模型

主机代码在CPU上运行并协调执行“设备”代码的“内核”。
设备代码执行在GPU和FGPA上，它是由内核建立的。
隐式的一套doall循环在：

工作组（work group）
工作组中的工作项（Work items）：迭代（尺寸通常为1-3）可以按照任意顺序执行，当被一个障碍物同步的时候，即遇到障碍（barrier）时，所有的工作项必须到达它（障碍物）为了超越它。

__kernel void mat_mul(const int N, __global float *A,
__global float *B, __global float *C
{
   	int i, j, k;
   	i = get_global_id(0);
   		j = get_global_id(1);
   			C[i*N+j] = 0.0f;
   			for (k = 0; k < N; k++) {
   				// C(i, j) = sum(over k) A(i,k) * B(k,j)
   				C[i*N+j] += A[i*N+k] * B[k*N+j];
   	}
}

可以看到OpenCL内核无外部循环，对工作项和工作组的隐式循环是在内核之外，并且可以以任意顺序执行。

工作项的多维领域

全局维度：10241024
局部维度：128128
可以选择一维，二维，三维中的一个用于你的算法
对于神经网络内的张量（tensors）而言，选三维

NDRange是一个N维索引空间，工作项也被映射其中，其中ND可以为1D，2D，3D.。
当主机提交一个内核给设备时会指定维度，NDRange和数据维度非常相近
Global_id = Group_id * Group_size + Local_id

OpenCL编程模型

三级分层执行结构匹配GPU结构
在这里插入图片描述
三级分层执行结构匹配CPU和FPGA结构

OpenCL存储器模型

三级存储器分层结构明确暴露给程序员：No attribute，local，global；如下图所示，寄存器单独给工作项使用，SRAM在同一工作组的工作项之间共享使用，DRAM是在工作组之间共享使用的。
存储器访问匹配计算并行性。
在这里插入图片描述
明确定义的内存层次结构：

全局数据：对应DRAM，共享于所有工作组之间
局部数据：对应SRAM，一个工作组中工作项之间共享
私密数据：对应registers，对应每个工作项

同行情况下，一个工作组中的工作项会获得连续的内存位置。

__kernel void k(__global float *A, …) {
__local Asub[M];
const int col = get_local_id(0);
Asub[col] = A[tCol*M+col];

context和命令队列（command queues）

context包括：

一个或多个设备
设备存储器
一个或多个命令队列
一个设备的所有命令（这些命令包括内核执行，同步化，和内存操作等）通过一个命令队列提交的。
每个队列指向context的一个设备。

内核之间的独立性

主机代码来管理内核之间的数据依赖性，控制依赖性包括：

顺序：就是一个接着一个
同步：同时全部内核发射，同时等待结束
流水线：一个刚发射，一个等待结束同时发射下一个时间。

内核完成的信号或者主机等待它或者另一个内核等待它。

矩阵乘法（matrix multiplication）

叠加（Tiling）内部循环来移动重复使用数据到片上SRAM。

SIMD机器的问题：conditonal

每一个条件下，线程被分成两个部分：一个执行条件为真的一支，然后执行条件为假的一支。
所以线程发散（divergence）是SIMD结构最大的性能损失问题。

GPU的总结！！！

GPU为整数和浮点数计算提供巨大的并行性
三级结构：一个设备包含多个SM（Stream Multiprocessors），一个SM包含多个核心。
运算模型匹配结构包含WG map to SM和WI map to cores 两种。

FPGA

Zynq UltraScale+ reconfigurable SOC

在这里插入图片描述

SOCs with ML accelerators

片上系统SOC

SOC是微处理器的一种极端进化，包含：

多CPU,GPU,VPs
存储器
特定应用加速器（通常是机器学习加速器）

Apple A14

在这里插入图片描述

HiSilicon Kirin 970

在这里插入图片描述

ASIC

Google Tensor Processing Unit

•ASIC devoted to matrix multiplications
• Third generation, TSMC 12nm
• 16K MAC/cycle, 2 GHz, 30 TFLOP/s BF16
• 250W, 120 TFLOP/J BF16 half-precision float
Each MXU (Matrix Multiplication Unit) has a systolic 256x256 architecture
– Activations flow along rows, weights along columns
– Output activations accumulate at bottom
在这里插入图片描述

Google Edge TPU

ASIC devoted to integer matrix multiplications
– Probably TSMC 12nm
• Supports 16-bit and 8-bit integer MACs
– Probably 64x64 systolic architecture
• Up to 4 TIOP/s, at 2W, about 2 TIOP/J
• Available on the Coral board

总结

VP，GPU，FPGA和ASIC，以下为总结表格：
在这里插入图片描述
• Datacenter platforms
– Deliver tens of TFLOP/s
– Routinely consume 100’s of W
– Deliver up to about 100 TFLOP/J
• Embedded platforms
– Deliver 100’s of GFLOP/s up to 1-2 TFLOP/s
– Consume up to 10 W (and down to 0.3 W)
– Deliver up to 500 TFLOP/J (only FPGAs)
• Energy consumption per operation is reduced via:
– Vectorized execution units (VPs and GPUs)
– Reconfigurable control and datapath (FPGAs)
– Custom ASIC accelerators (TPUs and some SOCs)
最后这些了解就好。