发掘 ARM GPU 的全部深度学习性能，TVM 优化带来高达2倍性能提升

雷锋网 AI 科技评论按：本文是由来自上海交通大学 Apex 实验室的本科生 Lianmin Zheng 发表于 TVM 的一篇博客，文中阐述了如何使用 TVM 优化移动端上的 ARM GPU 的深度学习。雷锋网 AI 科技评论对原文进行了编译。

随着深度学习取得了巨大成功，在移动设备上部署深度学习神经网络模型的需求也在迅速增长。与我们在桌面端平台所做的相类似，在移动设备上使用 GPU 可以同时实现加速推理计算和节约电能。但是现有的大多数深度学习框架并不能很好地支持移动端 GPU。问题的难点在于移动端 GPU 和桌面端 GPU 存在架构上的差异，这意味着需要投入更多专门的工作来实现移动端 GPU 的优化。正是这些额外的工作最终导致了大多数深度学习框架对移动端 GPU 的支持不足。

TVM 通过引入统一的 IR 栈来解决在不同硬件上的部署难题，通过这个 IR 栈可以轻松完成针对不同硬件的优化。在这篇文章中，我们展示了如何使用 TVM/NNVM 为 ARM Mali GPU 生成高效的内核，并进行端到端的编译（End-to-end compilation）。在我们基于 Mali-T860 MP4 的测试中，与 Arm Compute Library 相比，我们的方法在 VGG-16 上快了 1.4 倍，在 MobileNet 上快 2.2 倍。图形级别（Graph-level）和操作级别（Operator-level）的优化共同促进了这种加速。

在不同底层上测试 ImageNet 的推理速度

Mali Midgrad GPU

我们将使用带有 Mali-T860 MP4 的 Firefly-RK3399 作为我们的测试环境，所以我们下面主要关注 Mali T8xx。

架构

图 1 是 T860 和 T880 上的 Mali 架构图。GPU 可扩展到 16 个连通着色器核心（Coherent shader cores）。在每个着色器内核中，有 2 或 3 条运算流水线（Arithmetic pipelines），1 条加载/存储流水线（所谓的 TriPipe）。每个运算流水线中的 ALU 有四个 128 位向量单元和一个标量单元。我们使用 OpenCL 进行 GPU 计算。映射到 OpenCL 模型时，每个着色器核心负责执行一个或多个工作组。并且每个着色器核心最多支持 384 个并发执行的线程。OpenCL 中的每个工作项通常映射到 Mali GPU 上的单个线程。Mali GPU 使用 VLIW（超长指令字，Very Long Instruction Word）架构。每个指令字包含多个操作。Mali GPU 也可以使用 SIMD，因此大多数运算指令会在多个数据元素单元（Multiple data elements）上同时运行。^[1]

图1. Mali T860 和 T880（来源^[2]）

与英伟达 GPU 相比的不同点

与英伟达 GPU 相比，下面是我们在为 Mali GPU 编写 OpenCL 代码时需要关注的一些区别点。

• Mali GPU 使用统一的全局内存。在英伟达的 GPU 中，我们通常会将数据复制到共享内存中，因为英伟达的 GPU 在物理层面上将全局内存、共享内存和寄存器区分开了。在 Mali，这个复制操作并不会提高计算性能，因此可以移除这项操作。另外，Mali GPU 通常与 CPU 共享全局内存，所以 CPU 和 GPU 之间不需要数据的转移复制。

• Mali Midgrad GPU 是基于 SIMD（单指令多数据）而设计的，并且需要显性地进行向量化。在英伟达的 CUDA 中，并行性是通过 SIMT（单指令多线程）实现的，不需要显性地进行向量化。但是也要注意，较新的 Mali Bitfrost GPU 是基于四式矢量（Quad-style vectorization），并不需要显性地进行向量化。

• Mali GPU 中的所有线程都有独立的程序计数器。这意味着 warp 的大小为 1，所以分支发散（Branch divergence）不是一个大问题。

优化：以卷积操作为例

卷积层是大多数深度神经网络的核心，并且占用了大部分的计算时间。所以我们以卷积为例，说明如何在 TVM 中应用打包（Packing）、平铺（Tiling）、展开（Unrolling）和向量化（Vectorization）等常用技术。

使用 GEMM 实现 Im2Col

众所周知的卷积层算法是 im2col，它的原理是将小的 3D 输入立方体转换成矩阵的列并执行 GEMM 算法。这么做的优点在于，转化为矩阵运算之后可以使用高度优化的 BLAS 库。但是内存冗余问题（3x3 卷积存在 9 倍的内存冗余）也是相当可怕。

空间填充（Spatial Packing）

相反，我们采用另一种方法来计算卷积，并逐步应用一些优化技术。使用 VGG-16 中的卷积层作为微调样例，其配置如下所示。这里我们假设批量的大小为 1。

作为基准，我们还列出了 Arm Compute Library 中该层的性能。

声明计算过程：平铺和打包

平铺（Tiling）和打包（Packing）操作是用于更好地实现内存访问的两种方法。平铺操作将整个计算分成多个小块，以获得更好的数据重用（Data reuse）性能。包装操作则根据平铺重新排列输入矩阵，以便我们可以顺序地访问存储器，从而降低缓存未命中率。

我们在输入图像的宽度维度和滤波器矩阵的 CO 维度上进行平铺操作。这由代码 tvm.compute 进行声明。

# set tiling factor

VH = 1VW = 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）

下面的代码简单地将坐标轴（axes）绑定到 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：展开操作

循环展开（Loop 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 parame

ternum_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 con

v_, 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 来实现端到端编译。

性能

图2. 在不同底层上测试 ImageNet 的推理速度

如图2所示，我们在 ImageNet 上测试推理速度。在 Firefly-RK3399 上，Mali GPU 可以比 6 核 big.LITTLE 的 CPU 快 2 至 4 倍。我们的端到端流水线比 Arm Compute Library 快 1.4 至 2.2 倍。在 Arm Compute Library 中，我们尝试使用 GEMM 和直接卷积的方法，在这些测试用例中 GEMM 方法总是比直接方法快，所以我们只绘制了 GEMM 方法的结果。

图中缺失了一些结果，比如 Arm Compute Library 上的 resnet18，这是因为 Arm Compute Library 的图形运行时还暂时不支持跳转连接（Skip connection）操作，并且深度卷积（Depthwise convolution）的实现效果较差。这也反映了 NNVM 软件栈的优势。

半精度性能

深度神经网络的精度不是很重要，特别是对移动设备的推理过程而言。使用低精度算术可以使得推理速度更快。我们还测试了 Mali GPU 上的半精度浮点数。

表1. ImageNet 上 FP16 的推理速度

从理论上讲，FP16 既可以使得峰值计算加倍又可以使得内存开销减半，从而使速度提高一倍。但是对于较长的向量化和调优某些参数，它则需要更好的输入形状（Input shape）。

在移动设备上的更多工作

我们承认还有一些改进空间，它们主要是在图形层面。比如模型压缩和权重预布局。NNVM 的下一步改进将试图解决这些问题。

代码传送门

• End-to-End benchmark

• Convolution and Depthwise Convolution Schedule

引用

[1] ARM Mali GPU OpenCL Developer Guide

[2] ARM Developer

Via Optimizing Mobile Deep Learning on ARM GPU with TVM，由雷锋网 AI 科技评论编译。