TVM在ARM GPU上优化移动深度学习
随着深度学习的巨大成功，将深度神经网络部署到移动设备的需求正在迅速增长。与在台式机平台上所做的类似，在移动设备中使用GPU可以提高推理速度和能源效率。但是，大多数现有的深度学习框架都不能很好地支持移动GPU。困难在于移动GPU架构和台式机GPU架构之间的差异。这意味着在移动GPU上进行优化需要付出特殊的努力。繁琐的额外工作最终导致大多数深度学习框架中对移动GPU的支持不佳。
TVM通过引入统一的IR堆栈解决了部署不同硬件的困难，通过该IR堆栈可以轻松完成针对不同硬件的优化。本文展示了如何使用 TVM / NNVM为ARM Mali GPU生成有效的内核并进行端到端编译。在对Mali-T860 MP4的测试中，与Arm Compute Library相比，的方法在VGG-16上快1.4倍，在MobileNet上快2.2倍。图形级和算子级优化都有助于加快速度。

ImageNet上不同后端的推理速度图
MALI中级GPU
使用带有Mali-T860 MP4的Firefly-RK3399作为的测试环境，主要关注下面的Mali T8xx。
建筑学
图1是T860和T880上的Mali体系结构的概述。GPU最多可扩展到16个一致的着色器内核。在每个着色器内核内部，有2或3条算术管道，1条加载/存储管道和1条纹理管道（所谓的TriPipe）。每个算术流水线中的ALU具有四个128位向量单元和一个标量单元。
使用OpenCL进行GPU计算。映射到OpenCL模型时，每个着色器内核将执行一个或几个工作组。每个着色器内核最多支持384个并发执行的线程。OpenCL中的每个工作项通常都映射到Mali GPU上的单个线程。Mali GPU使用VLIW（超长指令字）架构。每个指令字包含多个操作。Mali GPU还使用SIMD，大多数算术指令可同时对多个数据元素进行操作。

图1. Mali T860和T880
与NVIDIA GPU的不同
与为NVIDIA GPU编写代码相比，在为Mali GPU编写OpenCL代码时，需要注意一些差异。
• Mali GPU使用统一的全局内存。在NVIDIA的GPU中，通常将数据复制到共享内存中，因为NVIDIA的GPU具有物理上独立的全局内存，共享内存和寄存器。在Mali，此副本不会提高性能，可以删除。此外，Mali GPU通常与CPU共享全局内存，无需在CPU和GPU之间进行复制。
• Mali Midgrad GPU基于SIMD（单指令多数据），并且需要显式矢量化。在NVIDIA CUDA中，并行性是通过SIMT（单指令多线程）实现的，而SIMT不需要显式矢量化。注意，较新的Mali Bitfrost GPU基于四边形矢量化，不需要显式矢量化。
• Mali GPU中的所有线程都有单独的程序计数器。这意味着warp size，因此分支分歧不是主要问题。
优化：以卷积为例
卷积层是大多数深度神经网络的内核，占用大部分计算时间。以卷积层为例来说明如何在TVM中应用诸如打包，平铺，展开和矢量化之类的常见优化技术。
Im2Col与GEMM
卷积层的一种著名算法是im2col，将小3D输入多维数据集转换为矩阵的列并执行GEMM。方法的优点是易于利用高度优化的BLAS库。内存冗余（3x3内核为9x内存）非常糟糕。
空间批处理
相反，采用一种方法来计算卷积，并逐步应用优化技术。VGG-16中的卷积层用作调整案例，其配置在下面列出。假设批处理大小为1以便进行推断。

作为基准，还在Arm Compute库中列出了该层的性能。
声明计算：平铺和打包
平铺和打包是旨在更好地访问内存的两种方法。平铺将整个计算分成小块，以实现更好的数据重用。打包根据平铺对输入矩阵进行重新布局，以便可以顺序访问内存，从而降低了缓存未命中率。
对输入图像的宽度尺寸和滤镜矩阵的CO尺寸进行平铺。通过tvm.compute来描述。

set tiling factor

VH = 1
VW = VC = 4

get input shape

_, CI, IH, IW = data.shape
CO, CI, KH, KW = kernel.shape
TH = IH + 2 * H_PAD
TW = IW + 2 * W_PAD

calc output shape

OH = (IH + 2H_PAD - KH) // H_STR + 1
OW = (IW + 2W_PAD - KW) // W_STR + 1

data shape after packing

dvshape = (N, TH // (VHH_STRIDE), TW // (VWW_STRIDE), CI, VHH_STRIDE+HCAT, VWW_STRIDE+WCAT)

kernel shape after packing

kvshape = (CO // VC, CI, KH, KW, VC)

ovshape = (N, CO // VC, OH // VH, OW // VW, VH, VW, VC)
oshape = (N, CO, OH, OW)

define packing

data_vec = tvm.compute(dvshape, lambda n, h, w, ci, vh, vw:
data_pad[n][ci][hVHH_STRIDE+vh][wVWW_STRIDE+vw], name=‘data_vec’)

kernel_vec = tvm.compute(kvshape, lambda co, ci, kh, kw, vc:
kernel[co*VC+vc][ci][kh][kw], name=‘kernel_vec’)

define convolution

ci = tvm.reduce_axis((0, CI), name=‘ci’)
kh = tvm.reduce_axis((0, KH), name=‘kh’)
kw = tvm.reduce_axis((0, KW), name=‘kw’)

conv = tvm.compute(ovshape, lambda n, co, h, w, vh, vw, vc:
tvm.sum(data_vec[n, h, w, ci, vhH_STRIDE+kh, vwW_STRIDE+kw].astype(out_dtype) *
kernel_vec[co, ci, kh, kw, vc].astype(out_dtype),
axis=[ci, kh, kw]), name=‘conv’)

unpack to correct layout

output = tvm.compute(oshape, lambda n, co, h, w:
conv[n][co//VC][h/VH][w//VW][h%VH][w%VW][co%VC],
name=‘output_unpack’, tag=‘direct_conv_output’)
通过以下方法检查定义的IR
print(tvm.lower(s, [data, kernel, output], simple_mode=True))
我在这里选择卷积部分。
produce conv {
for (co, 0, 64) {
for (h, 0, 56) {
for (w, 0, 14) {
for (vw.init, 0, 4) {
for (vc.init, 0, 4) {
conv[((((((((co*56) + h)*14) + w)*4) + vw.init)4) + vc.init)] = 0.000000f
}
}
for (ci, 0, 256) {
for (kh, 0, 3) {
for (kw, 0, 3) {
for (vw, 0, 4) {
for (vc, 0, 4) {
conv[((((((((co56) + h)*14) + w)*4) + vw)4) + vc)] = (conv[((((((((co56) + h)*14) + w)*4) + vw)4) + vc)] + (data_vec[(((((((((h14) + w)*256) + ci)*3) + kh)*6) + kw) + vw)]kernel_vec[((((((((co256) + ci)*3) + kh)*3) + kw)*4) + vc)]))
}
}
}
}
}
}
}
}
}
内核1：绑定线程
在TVM中，首先声明计算，然后调度。这种机制使算法和实现细节脱钩。（这个想法来自Halide）。
以下调度仅将轴绑定到GPU线程，代码可以在Mali GPU上运行。

helper function for binding thread

def tile_and_bind3d(s, tensor, z, y, x, z_factor=2, y_factor=None, x_factor=None):
“”" tile and bind 3d “”"
y_factor = y_factor or z_factor
x_factor = x_factor or y_factor
zo, zi = s[tensor].split(z, z_factor)
yo, yi = s[tensor].split(y, y_factor)
xo, xi = s[tensor].split(x, x_factor)
s[tensor].bind(zo, tvm.thread_axis(“blockIdx.z”))
s[tensor].bind(zi, tvm.thread_axis(“threadIdx.z”))
s[tensor].bind(yo, tvm.thread_axis(“blockIdx.y”))
s[tensor].bind(yi, tvm.thread_axis(“threadIdx.y”))
s[tensor].bind(xo, tvm.thread_axis(“blockIdx.x”))
s[tensor].bind(xi, tvm.thread_axis(“threadIdx.x”))

set tunable parameter

num_thread = 8

schedule data packing

_, h, w, ci, vh, vw = s[data_vec].op.axis
tile_and_bind3d(s, data_vec, h, w, ci, 1)

schedule kernel packing

co, ci, kh, kw, vc = s[kernel_vec].op.axis
tile_and_bind(s, kernel_vec, co, ci, 1)

schedule conv

_, c, h, w, vh, vw, vc = s[conv].op.axis
kc, kh, kw = s[conv].op.reduce_axis

s[conv].reorder(_, c, h, w, vh, kc, kh, kw, vw, vc)
tile_and_bind3d(s, conv, c, h, w, num_thread, 1, 1)

_, co, oh, ow = s[output].op.axis
tile_and_bind3d(s, output, co, oh, ow, num_thread, 1, 1)
有了这个时间表，的代码现在可以运行了，但是性能却很糟糕。

内核2：展开unrolling
循环展开可以减少循环控制的指令，减少分支惩罚并隐藏读取内存中的延迟。TVM通过调用以下命令轻松完成此操作s.unroll(axis)

set tunable parameter

num_thread = 8

schedule data packing

_, h, w, ci, vh, vw = s[data_vec].op.axis
tile_and_bind3d(s, data_vec, h, w, ci, 1)

“”"!! ADD UNROLL HERE !!"""
s[data_vec].unroll(vw)

schedule kernel packing

co, ci, kh, kw, vc = s[kernel_vec].op.axis
tile_and_bind(s, kernel_vec, co, ci, 1)

“”"!! ADD UNROLL HERE !!"""
s[kernel_vec].unroll(kh)
s[kernel_vec].unroll(kw)
s[kernel_vec].unroll(vc)

schedule conv

_, c, h, w, vh, vw, vc = s[conv].op.axis
kc, kh, kw = s[conv].op.reduce_axis

s[conv].reorder(_, c, h, w, vh, kc, kh, kw, vw, vc)
tile_and_bind3d(s, conv, c, h, w, num_thread, 1, 1)

“”"!! ADD UNROLL HERE !!"""
s[conv].unroll(kh)
s[conv].unroll(kw)
s[conv].unroll(vw)
s[conv].unroll(vc)

_, co, oh, ow = s[output].op.axis
tile_and_bind3d(s, output, co, oh, ow, num_thread, 1, 1)

内核3：矢量化
为了在Mali GPU上实现最佳性能，需要明确地进行矢量化。

set tunable parameter

num_thread = 8

schedule data packing

_, h, w, ci, vh, vw = s[data_vec].op.axis
tile_and_bind3d(s, data_vec, h, w, ci, 1)

unroll

s[data_vec].unroll(vw)

schedule kernel packing

co, ci, kh, kw, vc = s[kernel_vec].op.axis
tile_and_bind(s, kernel_vec, co, ci, 1)

unroll

s[kernel_vec].unroll(kh)
s[kernel_vec].unroll(kw)
“”"!! VECTORIZE HERE !!"""
s[kernel_vec].vectorize(vc)

schedule conv

_, c, h, w, vh, vw, vc = s[conv].op.axis
kc, kh, kw = s[conv].op.reduce_axis

s[conv].reorder(_, c, h, w, vh, kc, kh, kw, vw, vc)
tile_and_bind3d(s, conv, c, h, w, num_thread, 1, 1)

unroll

s[conv].unroll(kh)
s[conv].unroll(kw)
s[conv].unroll(vw)
“”"!! VECTORIZE HERE !!"""
s[conv].vectorize(vc)

_, co, oh, ow = s[output].op.axis
tile_and_bind3d(s, output, co, oh, ow, num_thread, 1, 1)

如何设置可调参数
至于上面的可调参数，可以计算一些。对于矢量化维VC，应该填充128位寄存器，因此对于float32可以将其设置为128/32 = 4，对于float16可以将其设置为128/16 = 8。
由于运行时间复杂，常常无法确定最佳值。在TVM中使用网格搜索。在TVM的高级IR中编写python代码，不是直接编写OpenCL代码，可以非常有效地完成。
生成的OpenCL代码
可以通过以下方式查看生成的OpenCL代码：
print(func.imported_modules[0].get_source())
OpenCL代码太长，无法粘贴到此处，由于展开太重而难以阅读。
端到端基准测试
比较一些流行的深度神经网络上不同后端之间的综合性能。测试环境是
Firefly-RK3399 4G
CPU: dual-core Cortex-A72 + quad-core Cortex-A53
GPU: Mali-T860MP4

Arm Compute Library : v17.12
MXNet: v1.0.1
Openblas: v0.2.18
使用NNVM和TVM进行端到端编译。
性能Performance

图2. ImageNet上不同后端的推理速度
如图2所示，在ImageNet上测试推理速度。在Firefly-RK3399上，Mali GPU的速度可以比6核big.LITTLE CPU快2倍至4倍。端到端管道比Arm Compute库快1.4倍至2.2倍。在Arm Compute Library中尝试了GEMM和卷积层的直接方法，在这些测试案例中，GEMM方法总是比直接方法快，因此仅绘制GEMM方法的结果。
图2中缺少一些结果，例如Arm Compute Library上的resnet18，因为Arm Compute Library的图形运行时当前不支持跳过连接，并且深度卷积的霓虹灯实现较差。这也反映了NNVM软件堆栈的优势。
半精度性能
深度神经网络的精度不是很重要，特别是对于移动设备的推断而言。使用低精度算术可以使推理更快。还在Mali GPU上测试了半精度浮点数。

从理论上讲，FP16既可以使峰值计算加倍，又可以使内存消耗减半，从而使速度加倍。需要良好的输入形状，以实现更长的矢量化和微调一些参数。
在移动设备上的进一步工作
还有一些改进的空间，主要是在图形级别，例如模型压缩和权重布局。NNVM的进一步改进将尝试解决这些问题。

TVM在ARM GPU上优化移动深度学习相关推荐

TVM 优化 ARM GPU 上的移动深度学习
TVM 优化 ARM GPU 上的移动深度学习随着深度学习的巨大成功,将深度神经网络部署到移动设备的需求正在迅速增长.与桌面平台上所做的类似,在移动设备中使用 GPU 既有利于推理速度,也有利于能源 ...
CUDA上的量化深度学习模型的自动化优化
CUDA上的量化深度学习模型的自动化优化深度学习已成功应用于各种任务.在诸如自动驾驶汽车推理之类的实时场景中,模型的推理速度至关重要.网络量化是加速深度学习模型的有效方法.在量化模型中,数据和模型参 ...
如何在 GPU 上优化卷积
如何在 GPU 上优化卷积将演示如何在 TVM 中编写高性能卷积实现.正方形大小的输入张量和过滤器为例,假设卷积的输入具有大batch批量.在这个例子中,使用不同的布局存储数据,实现更好的数据局部性 ...
如何在GPU上优化卷积
如何在GPU上优化卷积本文将演示如何在TVM中编写高性能的卷积实现.以平方大小的输入张量和滤波器为例,并假设卷积的输入量很大.使用不同的布局来存储数据,以实现更好的数据局部性.缓冲区布局为HWCN, ...
【Paper Reading】MCUNet：物联网设备上的微型深度学习
MCUNet: Tiny Deep Learning on IoT Devices 本文纯翻译,原文见:https://arxiv.org/abs/2007.10319 目录 MCUNet: Tiny ...
02.改善深层神经网络：超参数调试、正则化以及优化 W1.深度学习的实践层面（作业：初始化+正则化+梯度检验）
文章目录作业1:初始化 1. 神经网络模型 2. 使用 0 初始化 3. 随机初始化 4. He 初始化作业2:正则化 1. 无正则化模型 2. L2 正则化 3. DropOut 正则化 3.1 ...
02.改善深层神经网络：超参数调试、正则化以及优化 W1.深度学习的实践层面
文章目录 1. 训练,验证,测试集 2. 偏差,方差 3. 机器学习基础 4. 正则化 5. 为什么正则化预防过拟合 6. dropout(随机失活)正则化 7. 理解 dropout 8. 其他正则 ...
supervessel-免费云镜像︱GPU加速的Caffe深度学习开发环境
开发环境介绍在SuperVessel云上,我们为大家免费提供当前火热的caffe深度学习开发环境.SuperVessel的Caffe有如下优点: 1) 免去了繁琐的Caffe环境的安装配置,即申请即 ...
在腾讯云主机上搭建pytorch深度学习环境
在腾讯云主机上搭建pytorch深度学习环境 1.购置[腾讯云主机](https://cloud.tencent.com/? fromSource=gwzcw.234976.234976.234976 ...

TVM在ARM GPU上优化移动深度学习

set tiling factor

get input shape

calc output shape

data shape after packing

kernel shape after packing

define packing

define convolution

unpack to correct layout

helper function for binding thread

set tunable parameter

schedule data packing

schedule kernel packing

schedule conv

set tunable parameter

schedule data packing

schedule kernel packing

schedule conv

set tunable parameter

schedule data packing

unroll

schedule kernel packing

unroll

schedule conv

unroll

TVM在ARM GPU上优化移动深度学习相关推荐

最新文章

热门文章