分享嘉宾：邱东洋 腾讯 高级开发工程师
编辑整理：张智为
出品平台：DataFunTalk

导读： 随着业务规模的不断发展，算法模型复杂度不断增加，实时性要求很高的场景，对在线推理优化提出很大挑战。本文将和大家分享腾讯智能对话产品中模型推理优化的常见方法和聚焦GPU推理的方法论。主要内容包括以下几大方面：

• 背景介绍

• 推理性能优化的常用方法

• GPU并行加速的方法论

• 总结

  01 背景介绍

腾讯游戏知几是腾讯游戏的智能对话机器人，这款产品主要功能包括QA对话、闲聊、语音陪伴助手等，已覆盖和平精英、王者荣耀、天涯明月刀手游、QQ飞车等100+腾讯游戏。

从行业来看，基于transformer预训练模型出现以后，尤其是bert模型，各种NLP任务的效果都有了质的飞跃。然而这些模型越来越复杂，参数量非常的大，导致在使用这些模型的时候存在较大的性能瓶颈，实现在线实时推理几乎成为一件不可能的事情，所以做好推理加速是一个非常棘手的技术问题。同时，在游戏流量非常大的场景下，在性能上做好优化可以为在线服务带来质的飞跃，提高机器负载会使成本极大降低。

02 推理性能优化的常用方法

推理性能优化的常用方法主要包括两大类，一类是模型的压缩技术，比如模型的剪枝、量化、蒸馏等，在较大的预训练模型下被广泛使用；另外是推理加速的技术，在CPU、GPU加速上分别有一些方法。

1. 网络剪枝

在神经网络模型中，存在大量冗余的神经元和权重，它们对最终的结果产生的影响的权重又很低。基于此，我们如果对大的模型采取一些有效的压缩手段，就可以在保证模型效果的前提下，让模型变得更小，方便在线部署。

网络剪枝是从大型网络中筛选出不重要的神经元以及权重，将它们从网络中删除，同时尽可能地保留网络的性能。

剪枝主要分为两大类，如上图所示：一个是非结构化的剪枝，一个是结构化的剪枝。

• 结构化剪枝

剪除的基本单元是神经元，这种方式在现有的硬件条件下，可以实现比较明显的推理加速和存储的优势。缺点是剪枝力度有点大，会对模型精度产生较大的影响，导致模型效果下跌，无法达到预期效果。

• 非结构化剪枝

剪除的是单个权重，这样对精度的损失就会小一些，但最终产生的是稀疏矩阵。这种情况下就需要下层的硬件以及计算库有良好的支持，才能实现最终的推理加速和存储优势。

稀疏动量剪枝，是从头训练的一种方法，它的目的是在给定的网络稀疏度下，选择更有效的连接方式。

算法首先是要确定网络的剪枝比例，作为网络的稀疏度，构建初始的稀疏网络，然后再开始训练初始网络。初始网络不是最优的结构，要确认哪些连接需要删除，哪些需要重新添加网络。通过这种先剪枝后生长的方式可以达到更优的目标。

2. 量化技术

量化技术是把高精度表示的网络权重和激活值，用低精度来近似表示，实现网络的轻量化。优势如下:

• 网络存储 ：每个层权重量化后，32位的比特就可以压缩到8比特，就是浮点型到整形的量化，整个模型占的空间就会变小；
• 激活值 : 通过使用较少位的数值表示，在处理同样数据时需要读/写的内容就更短，内存带宽的压力就变得更小；
• 计算时间 ：单位时间内处理定点运算指令就会比浮点运算的指令多。

量化一般分为对称和非对称。

对称量化是通过收缩因子，将32的浮点型的最大值映射到8比特的最大值，这是一种线性的量化技术。

如果原始的这个浮点型分布不均匀的话，就会造成大量的数据集中。这是一种不饱和的线性量化，数据扎堆在某一个范围内。这样的话会存在一些范围利用少的情况，导致精度的损失比较大。

为了解决非饱和截取带来的区间浪费以及数据失真问题，英伟达提出饱和截取，即截取一段数据分布较为均匀的区域映射到-127至127之间。具体操作是计算一个阈值T，将-T～T之间的区域进行线性映射，而超出区域的数据点则统一赋予边界值+-127。这样做可以保证数据的大部分分布跟之前是一样的，同时将超出截断值的数据归一化到边界上，避免了数据的丢失。

尝试不同的T值，然后返回使量化后分布与原分布差距最小的T值。可以用分bins的方法来得到一组数据的分布，再利用KL距离来衡量量化前后的数据分布是否相似。

3. 模型蒸馏

蒸馏是一种模型压缩常见方法，将复杂、学习能力强的网络学到的特征，表示“知识”蒸馏出来，传递给参数量小、学习能力弱的网络。模型的参数量和所能捕获的知识量之间，其实并非一个稳定的线性关系，如下图所示：

并非如图中曲线1，而是接近边际效益，边际收益逐渐减少的一种增长曲线，如曲线2、3。这两个曲线还说明完全相同模型架构和模型的参数量，使用完全的训练数据，能获得的知识量也并不一定完全相同，这主要是训练方法的不同。

知识蒸馏需要两种类型的网络：Teacher Model和Student Model。前者参数量大、结构复杂，后者参数量较小、结构相对简单。二者可以是不同的网络结构，但是采用相似的网络结构，蒸馏效果会更好。训练流程如下图所示：

首先训练Teacher Model，然后用其指导Student Model的训练。将Teacher Model在Softmax层的输出作为数据的soft label，Student Model的loss function将是对soft label预测和hard label预测的loss的线性加权和。Student Model训练好后，按照常规模型使用即可。知识蒸馏通过将Teacher Model的知识迁移到Student Model中，使Student Model达到与Teacher Model相当的性能，同时又能起到模型压缩的目的。其局限性在于，由于使用Softmax层的输出作为知识，所以一般多用于具有Softmax层面的分类任务，知识蒸馏在其它任务上的表现并不好。

4. Caching

CPU频率远快于主存访问速度，在处理器时钟周期内，CPU常常需要等待主存，浪费计算资源。为了缓解CPU和内存之间速度的不匹配问题，增加了CPU cache 来解决。

cache 利用局部性原理来提高缓存命中率:

• A. 时间局部性：如果某个数据被访问，那么在不久的将来它很可能被再次访问；
• B. 空间局部性：如果某个数据被访问，那么与它相邻的数据很快也可能被访问。

代码匹配好局部性规则，可以提高缓存命中率进而提升性能。下面举一个矩阵乘法的实例，如下图所示:

上面遍历计算符合矩阵计算公式，行列相乘得到结果。数组在内存中是按照行进行存储的。上面计算的内存读取流程如下图所示，缓存矩阵B的缓存命中率很低。

交换j,k遍历顺序，计算结果不变，内存读取示意如下图所示，矩阵B按行读取，缓存命中率显著提高。

5. 指令集加速

英特尔指令集扩展是可提高性能的附加指令。当同一操作在多个数据对象上执行时，SIMD即单指令多数据的一种方式，是向量计算的一种方法，可以提高计算吞吐量。对于浮点数向量计算，用标量和向量指令集实现方式对比如下：

当向量很大时，标量版本循环次数很多，效率低。指令集细节可以查阅英特尔官网。

6. 多算子融合

算子融合是GPU上一个很重要的推理加速的一个优化手段，尤其是针对NLP这样的大模型，会带来比较显著的效果的提升。对于GPU异构编程，每一次op操作都会有一个内核的调用和多次的显存的读取；对于小op来说启动GPU kernel的时间会大于GPU计算时间，显存的读取开销也很大；op数目太多的话,效率会变低；所以将算子合并，可以有效地提高计算的性能。

对照下图AlexNet网络结构，我们可以进行垂直和水平的算子融合；Conv + Bias + ReLU 垂直融合为CBR算子；水平方向将输入为相同张量且执行相同操作的层融合在一起。

还有Parallel Reduction，Remove Padding等优化手段，这里就不展开说了。

03 GPU并行加速的方法论

1. GPU概述

GPU本来的任务是做图形图像的，图像有个特点就是并行度很高，属于并行任务。随着深度学习的发展，大家逐渐将目光转向GPU并行计算。x86 CPU+GPU的这种异构应该是最常见的，通过异构计算来提高吞吐量。下图是异构框架示意图:

左图：一个四核CPU一般有四个ALU，ALU是完成逻辑计算的核心，也是我们平时说四核八核的核，控制单元，缓存也在片上，DRAM是内存不在片上，CPU通过总线访问内存，CPU适合低并行逻辑复杂程序。

右图：GPU，绿色小方块是ALU，我们注意红色框内的部分SM，这一组ALU公用一个Control单元和Cache，这个部分相当于一个完整的多核CPU，但是不同的是多了很多计算核心，而control部分变小，可见计算能力提升了，控制能力减弱了；GPU适合高并行逻辑简单的大数据计算。

CPU/GPU线程区别：

• CPU线程是重量级实体，操作系统交替执行线程上下文切换花销很大；GPU线程是轻量级的，包括成千上万个，多数在排队状态，线程之间的切换基本没有开销。
• CPU的核被设计用来尽可能减少一个或二个线程运行时间的延迟；而GPU则是大量线程，最大幅度地提高吞吐量。

2. CUDA基础

(1) 编程结构

一个完整的CUDA应用可能的执行顺序如下图：

从host的串行到调用核函数-->核函数被调用后控制马上归还主机线程-->继续执行host代码。

(2) 内存管理

• 主机和设备之间的数据传输
  设备与GPU之间的峰值理论带宽远高于主机和设备之间的理论带宽；
  基于这种设计，我们在使用时应尽量避免主机和设备之间的数据传输；
  优化：固定内存、异步、统一虚拟地址等。
• 设备上有层次的内存空间
  全局内存：资源丰富，访问延时最大；
  共享内存：在片上，访问更快。

(3) 线程管理

• 分层结构让并行过程更佳灵活
  如上图所示，1个核函数调用只能有1个网格，但1个Grid可以包含多个块，1个块又可以包含多个线程。在网格中块的组织形式，以及在块中线程的组织形式，可以是一维、二维或三维，这样就可以在并行计算中有很灵活的组织方式。
• 块内/块间
  块内可以实现线程同步，共享内存；
  块间物理隔离，不能相互影响。
• 线程编号
  块的维度和块的索引相乘，再加上线程的索引。

(4) 执行模型

下图从逻辑角度和硬件角度描述CUDA编程模型对应的组件。

SM中共享内存和寄存器是关键的资源，线程块中线程通过共享内存和寄存器相互通信协调。因为SM有限，虽然我们的编程模型层面看所有线程都是并行执行的，但是在微观上看，所有线程块也是分批次地在物理层面的机器上执行，线程块里不同的线程可能进度都不一样，但是同一个线程束内的线程拥有相同的进度。

并行就会引起竞争，多线程以未定义的顺序访问同一个数据，会导致不可预测的行为，CUDA只提供了一种块内同步的方式，块之间没办法同步。

同一个SM上可以有不止一个常驻的线程束，有些在执行，有些在等待，他们之间状态的转换是不需要开销的。

3. 性能分析工具

在使用tensorflow的过程中，我们经常需要使用工具来监测模型的运行性能。合理的利用优化工具可以提高工作效率，让我们把主要精力花在关键代码的优化上。

nvvp，nvprof是cuda toolkit集成的工具，用于生成GPU timeline的工具。nvprof是命令行工具，我们的模型常常是运行在远端的服务器上，我们需要把输出的监测数据拷贝至本地查看，这个时候需要用到nvvp进行可视化分析。nsight是NVIDIA最新的用于监测kernel timeline的工具。

下图是nvvp可视化分析各个函数在调用过程中的时间占比情况。

上图是一个矩阵加法例子，可以看到通过nvprof工具可以分析到GPU活跃活动中最耗时的部分是数据从主机到设备的传输，它占了64%的时间。还可以看到API调用的一些热点的分析，CudaMalloc这个函数被调用了三次，耗时大概500ms，占比在一半以上，通过nvprof命令可以看到整个程序执行热点的概述。

还可以通过这个指令结合参数分析程序具体的一些指标，分析它的achieved_occupancy（占用率），gld_throughput（内存吞吐）和gld_efficiency（内存使用效率）。占用率是在sm上活跃的线程数，占最大的线程数的比例。

可以看到下图中是0.7，这已经非常高。除了这三个指标之外，其实还有很多指标可以使用，具体可以参考NVIDA官网。

4. 理论计算极限

我们在优化之前需要搞清楚一个问题，我们这次优化究竟可以取得多大收益？

这里介绍两个定律，可以帮助大家在优化过程中做一个简单理论上的评估。

先引入强扩展定义：规模固定的情况下，通过提高多处理器的数量来减少代码执行的时间。

对应Amdahl's Law(阿姆达尔定律)：提升一个系统的一个部分的性能对整个系统有多大影响。问题转变为怎么测算加速比S，P是可以并行化代码分数，N是并行部分可以得到的多处理器的数量。通过公式我们可以看到P越大，加速比就会越大，意味着优化空间就越大。如果P很小的情况下，不论增加多少计算资源，其实整个收益都会不佳。

对应弱扩展：问题的规模可能会随着资源的增加而增加，可以用古斯塔夫森定律来理论分析，这里加速比可以看随着N的增加，整个问题规模也是越来越大。我们在优化的时候首先就是要分析问题属于哪一类场景，再根据每类场景对应公式做一个简单的推测。

在优化过程当中，每一项指标其实都是受到硬件的限制，那么每个指标的理论的峰值天花板是多少呢？这里也给大家简单介绍一下。

Tesla K10 单精度FLOPS极限：

• 745 MHz 内核时钟 * 2 GPU * (8 多处理器 * 192 浮点计算/ 多处理器) * 2 ops/cycle = 4.58 TFLOPS

Tesla K10 内存带宽极限：

• 2 GPU * 256 bit * 2500 MHz 内存时钟 * 2 DDR / 8 bits/ byte =320 GB/s

指令比 bytes：

• 4.58 TFLOPS / 320 GB/s yields 14.3 instructions: 1 byte

5. APOD方法论

• Assess ：对于一个项目，第一步要做的是接触 (Assess) 项目，得到项目代码中每部分的执行时间。有了这部分内容，开发者就可以找到并行优化的瓶颈所在，并开始尝试用GPU加速。根据前面提到的Amdahl's and Gustafson's laws，可以确定并行优化的性能上界。
• Parallelize ：找到瓶颈所在并确定了优化的目标和期望，就可以优化代码。调用一些如cuBLAS, cuFFT, orThrust的GPU-optimized library可能会很有效。另一方面，有些应用需要开发者重构代码，以此让可以被并行优化得部分被暴露出来。
• Optimize ：确定了需要被并行优化的部分之后，就要考虑具体的实现方式了。具体的实现方式通常只有一种，所以充分理解应用的需求是很有必要的。要记得，APOD 是一个反复迭代的过程（找到可以优化的点，实现并测试优化，验证优化效果，然后再重复）。因为对于开发者来说，没有必要最初就找到解决所有性能瓶颈的策略。优化可以在不同的 level 上进行，配合性能分析工具是很有帮助的。
• Deploy ：大的原则是当优化完一处之后，立刻将这一部分部署到生产环境，而不是再去寻找其他可以优化的地方。这样做有很多重要的原因，比如这会使得用户尽早从这个优化中收益。逐步迭代的方式，有助于减少风险，保证线上的稳定性。

6. 优化案例分享

(1) 矩阵加法

对比CPU，GPU执行矩阵加法，对比耗时变化；通过nvprof分析GPU执行性能，可以看到GPU的整体在GPU上虽然执行核函数只有十毫秒，但是整体的耗时也差不多有800ms，主要是从设备到储存的数据传输，耗时比较大。代码可以重点看第四步核函数的实现部分，可以改变grid，block来看性能的变化。

(2) 矩阵乘法

该例子分析使用共享内存和不使用共享内存性能差距。在Tesla P40 GPU机器上，对于16000多维的矩阵乘法矩阵，不使用共享内存耗时40多秒，而使用共享内存时间大概是8秒多，下图展示了使用共享内存减少对全局内存的访问次数分析。

使用共享内存设计核函数，下面两个概念很重要：

• 跨内存存储体映射数据元素
• 从线程索引到共享内存偏移的映射

GPU共享内存并不是线性的，而是二维的，分成不同存储体的，并行也不是循环。对于一个二维的块，数据是按照行放进存储体中，线程束中取数据按照行来进行的访问方式是最优x[threadIdx.y][threadIdx.x]，threadIdx.x在线程束中为连续变化的，对应到共享内存中也是遍历共享内存的同一行的不同列。

04 总结

最后总结一下今天的分享内容。首先介绍了推理优化的“术”，即常见方法，包括模型压缩（剪枝、量化和整理）、推理加速（Caching、指令集加速、多算子融合、其他）等；后面介绍了推理加速的“道”，聚焦GPU讲了一些方法论，包括理解硬件，分析问题，建立目标和APOD循环迭代，还介绍了问题分析工具以及理论峰值的预估。本文介绍的是相对基础的知识，大家实践过程当中可以参考CUDA官网来深入研究。

分享嘉宾：