思享工具箱

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

摘要：
TVM优化了ARMGPU上的移动深度学习随着深度学习的巨大成功，将深度神经网络部署到移动设备的需求正在快速增长。这意味着移动GPU的优化需要特别的努力。TVM通过引入统一的IR堆栈解决了部署不同硬件的困难，该堆栈可以轻松完成对不同硬件的优化。本文介绍了如何使用TVM/NNVM为ARMMaliGPU生成有效内核并对其进行端到端编译。OpenCL中的每个工作项通常映射到MaliGPU上的单个线程。MaliGPU使用统一的全局内存。MaliMidgradGPU基于SIMD，需要显式矢量化。
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倍。图形级和算子级优化都有助于加快速度。

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

 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，大多数算术指令可同时对多个数据元素进行操作

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

 图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以便进行推断。

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

 作为基准，还在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 + 2*H_PAD - KH) // H_STR + 1
OW = (IW + 2*W_PAD - KW) // W_STR + 1
 
# data shape after packing
dvshape = (N, TH // (VH*H_STRIDE), TW // (VW*W_STRIDE), CI, VH*H_STRIDE+HCAT, VW*W_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][h*VH*H_STRIDE+vh][w*VW*W_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, vh*H_STRIDE+kh, vw*W_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[((((((((co*56) + h)*14) + w)*4) + vw)*4) + vc)] = (conv[((((((((co*56) + h)*14) + w)*4) + vw)*4) + vc)] + (data_vec[(((((((((h*14) + w)*256) + ci)*3) + kh)*6) + kw) + vw)]*kernel_vec[((((((((co*256) + 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)

有了这个时间表，的代码现在可以运行了，但是性能却很糟糕。

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

内核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)

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

内核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)

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

如何设置可调参数

至于上面的可调参数，可以计算一些。对于矢量化维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

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

 图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上测试了半精度浮点数。

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

 从理论上讲，FP16既可以使峰值计算加倍，又可以使内存消耗减半，从而使速度加倍。需要良好的输入形状，以实现更长的矢量化和微调一些参数。

在移动设备上的进一步工作

还有一些改进的空间，主要是在图形级别，例如模型压缩和权重布局。NNVM的进一步改进将尝试解决这些问题。

 

标签：#axis#kw#conv#ci | 浏览：116 | 发布日期：

免责声明：文章转载自《TVM在ARM GPU上优化移动深度学习》仅用于学习参考。如对内容有疑问，请及时联系本站处理。

上篇C#实现WinForm随WINDOWS服务一起启动(转载)测试人员需要了解的sql知识（基础篇）下篇

宿迁高防，2C2G15M，22元/月；香港BGP，2C5G5M，25元/月 雨云优惠码：MjYwNzM=

相关文章

手把手教你用Jenkins CI 自动部署Docker + 使用阿里云镜像服务自动构建+ webhook触发

Jenkins部分 首先,我们要有个Jenkins咯,下载链接:https://jenkins.io/download/ 我们安装官网教程安装好jenkins,安装教程略.... 嗯?不是说好手把手么?你妹的. 好好好,我们还是来手把手教程好了. 首先安装JDK8 添加安装源之后直接apt-get install就好,下面是ubuntu的安装命令,其他系统...

Java安全之Axis漏洞分析

Java安全之Axis漏洞分析 0x00 前言 看到个别代码常出现里面有一些Axis组件，没去仔细研究过该漏洞。研究记录一下。 0x01 漏洞复现 漏洞版本：axis=<1.4 Axis1.4 freemarker 下载Axis包1.4版本将Axis放到tomcat的webapp目录中。freemarker.jar放到Axis的 lib目录下。运行t...

目标检测相关基础概念

目标检测相关基础概念 算法分类 flowchart LR 1[Object Detection] 2[two stage]-->R-CNN-->SPP-NET-->A[Fast R-CNN]-->B[Faster R-CNN] 3[one stage]-->a[OverFeat] 1-->2 1-->3 a--&g...

如何在 Matlab 中绘制带箭头的坐标系

如何在 Matlab 中绘制带箭头的坐标系 如何在 Matlab 中绘制带箭头的坐标系 实现原理 演示效果 完整代码 实现原理 使用 matlab 的绘制函数时，默认设置为一个方框形的坐标系， 图1 如果想要绘制的如下图所示中的带箭头的坐标系，需要如何实现呢？ 图2 方法一：通过设置 axis 对象 的属性来调整坐标轴，参考代码如下： % 通过...

微信小程序访问webservice（wsdl）+ axis2发布服务端(Java)

0、主要思路：使用axis2发布webservice服务端，微信小程序作为客户端访问。步骤如下： 1、服务端： 首先微信小程序仅支持访问https的url，且必须是已备案域名。因此前期的服务器端工作需要先做好，本人是申请了个人域名（已备案），并使用阿里云服务器，然后申请免费SSL，通过配置tomcat完成支持https访问。此外，intellJ IDE的j...

大型前端项目 DevOps 沉思录 —— CI 篇

本文作者：成龙 腾讯前端开发工程师，负责腾讯文档前端开发与研发效能提升，AlloyTeam成员。 导语 本篇文章将着重探讨 DevOps 在持续集成阶段需要提供的能力，将对工作流的设计及流水线的优化思路做一个简要讲解。 DevOps 一词源于 Development 和 Operations 的组合，即将软件交付过程中开发与测试运维的环节通过工具链打通，并...

最新文章

随机推荐

思享工具箱导航