Chiphell - 分享与交流用户体验

 找回密码
 加入我们
搜索
      
查看: 10493|回复: 16

新Titan X的INT8计算到底是什么鬼

[复制链接]
发表于 2016-7-23 13:26 | 显示全部楼层 |阅读模式
本帖最后由 我輩樹である 于 2016-7-23 13:36 编辑

新Titan X发布,兴奋的睡不着,起来刷paper,目前能找到的资料不多,多来自于google。如下:
1,On the efficient representation and execution of deep acoustic models
http://arxiv.org/abs/1607.04683
介绍如何开启使用INT8的大门

2,Fast, Compact, and High Quality LSTM-RNN Based Statistical Parametric Speech Synthesizers for Mobile Devices
http://arxiv.org/abs/1606.06061
介绍了一个基于RNN使用INT8的一个语音识别的例子

在传统思维里面,似乎科学计算应该是精度越高越好,然而这几年老黄推波助澜,精度别说fp64了,fp32,fp16,直到今天,浮点都有被抬出场外的风险(当然仅限于深度学习)。
为什么说深度学习这么独特,居然不需要多高精度呢,主要原因还是在于其统计学算法的本质。这种算法本身就是一种纠错过程,试图纠正模型预测与现实的差距,学名“最小化代价函数”。精度的损失,也可以一同算入和现实的差距中,由算法过程来自行减少。
当然是否真的有那么神奇,因为时间关系我并没有去验证,虽然理论上还有种种需要细心处理的部分,但至少应该可行的。
那么新Titan X高达44 T flops的INT8是否能成全老黄新一代的霸主地位呢?

就目前的资料来看,INT8精度主要用在模型的inference上。
无论训练的多么精妙的模型,最终还要给用户使用才行。这个使用的过程即所谓的模型infernece。
简单不恰当的说,拿汽车做例子,模型的训练就是制造变速箱的过程,inference就是变速箱实际运作挂档的过程。

INT8用于训练流程,仍然是不太足够的,至少我还没找到相关paper。主要是做模型的使用,特别是在手机设备上。

google把它用于语音识别,是因为它迫切需要降低运算和储存/传输的压力。如果把经典深度神经网络(如cnn)看做一张网,那么用于语音识别的RNN就是一堆网叠加在一块,彼此互联。因为语音是一个时序数据,每一个时间点的信号都需要一整张网来模拟。可想而知数据流之大。
论文中也说明了可以显著的减少储存和运算的需求。

INT8的FMA操作。
如果把神经网络分解到最后,绝大部分都会对应FMA操作。
FMA操作:
Screenshot from 2016-07-23 13:10:10.png

很简单吧,其实没啥复杂的,只不过里面的大字母都代表矩阵。
Screenshot from 2016-07-23 13:12:08.png


这个处理中,与普通的fma多了两个流程。
X(float)经过Q函数处理成uint8。这就是google所谓的线性量化过程,将float的数值拉皮成到int8范围内。
同时W(uint8)直接就int8了,这俩相乘,生成一个int32的中间变量。
注意这个Mult函数就是整个深度学习中使用最频繁,计算占最大比重的矩阵乘法。而且多出来的量化与反量化过程同样可以通过simd加速,对性能的影响微乎其微。(The overhead of the quantization and recovery operations
is typically negligible, and also parallelizable via SIMD)
也就是说最大比重的指令将在两个int8输入下完成,足以说明一点:
老黄没有吹牛,基于int8的新深度学习构架很快就会到来。(google已帮你免费验证了)

当然,拉完皮后一定会损失精度,何况在后续的算法中,还要重新拉回float参与后续计算(途中的R函数)。如何避免损失,google引入了一个新的梯度下降法Quantization aware training,有兴趣的可以看论文。

发表于 2016-7-23 13:43 | 显示全部楼层
楼主学CS的吧。。。
 楼主| 发表于 2016-7-23 13:44 | 显示全部楼层
wujin941005 发表于 2016-7-23 13:43
楼主学CS的吧。。。

cs博狗。
发表于 2016-7-23 13:45 | 显示全部楼层
 楼主| 发表于 2016-7-23 13:46 | 显示全部楼层

别别别,还没拿到那张纸呢,差别老大了。
发表于 2016-7-23 14:16 来自手机 | 显示全部楼层
本帖最后由 fairness 于 2016-7-23 14:42 编辑
我輩樹である 发表于 2016-7-23 13:44
cs博狗。

膜拜,LZ读博期间能否搞个接地气的项目: 基于深度学习的A股投资策略,主要的training过程在刚刚诞生的太湖之光超算上完成。绝对名利双收!
发表于 2016-7-23 14:18 | 显示全部楼层
一个周期能执行8次int8运算,应该要增加额外的硬件电路吧
发表于 2016-7-23 14:41 | 显示全部楼层
本帖最后由 chungexcy 于 2016-7-23 15:07 编辑

还是觉得哪点有问题。

Google那个图,要求w*x的结果是uint32。但按照FMA c=a*b+c,一般情况下,c应该和a、b类型相同。而向量乘法需要很多个fma,即使最后的结果可以uint32,中间的c不太可能能在uint8以内。

除非nvidia有黑科技,c能保证uint32的精度的特殊fma整数操作。按照带宽来说,这个我是不信的。
不对,c可以一直缓存着,好像也是有可能的,a, b是uint8,c是uint32,但一个cuda能放下4个uint32么==


比如,intel 的 AVX512IFMA52 指令集:
__m512i _mm512_madd52lo_epu64 (__m512i a, __m512i b, __m512i c)

Synopsis
  1. __m512i _mm512_madd52lo_epu64 (__m512i a, __m512i b, __m512i c)
  2. #include "immintrin.h"
  3. Instruction: vpmadd52luq
  4. CPUID Flags: AVX512IFMA52
复制代码

Description
  Multiply packed unsigned 52-bit integers in each 64-bit element of b and c to form a 104-bit intermediate result. Add the low 52-bit unsigned integer from the intermediate result with the corresponding unsigned 64-bit integer in a, and store the results in dst.

Operation
  1. FOR j := 0 to 7
  2.             i := j*64
  3.             tmp[127:0] := ZeroExtend64(b[i+51:i]) * ZeroExtend64(c[i+51:i])
  4.             dst[i+63:i] := a[i+63:i] + ZeroExtend64(tmp[51:0])
  5. ENDFOR
  6. dst[MAX:512] := 0
复制代码

Intel专门分了两个指令来存储高52和低52。就算是这个,要实现一个整数fma,也是两条指令了。这样算下来还不如fp16。
__m512i _mm512_madd52lo_epu64 (__m512i a, __m512i b, __m512i c)
__m512i _mm512_madd52hi_epu64 (__m512i a, __m512i b, __m512i c)

发表于 2016-7-23 14:55 | 显示全部楼层
PolyMorph 发表于 2016-7-23 14:18
一个周期能执行8次int8运算,应该要增加额外的硬件电路吧

估计是simd设计,uint8刚好是fp32的1/4大小,一个cuda算4个uint8操作。
发表于 2016-7-23 15:07 | 显示全部楼层
这是啪啪打google那个TPU的脸吧。之前google说一个order of magnitude的能耗比提升。老黄pascal x2倍int8 x4倍,直接迎头一棒。

具体提升inference速度和能耗比有很多思路,变精度是比较实际的一种。长期我更看好binary neuron。
发表于 2016-7-23 19:16 | 显示全部楼层
我只能问,你是打CS1.5还是CS1.6……

帮顶吧,titan这HBM2,没戏了?
发表于 2016-7-23 19:18 | 显示全部楼层
CS是什么,computer science?
发表于 2016-7-23 20:57 | 显示全部楼层
如果单位今年不扣我钱 我想买个
发表于 2016-7-23 22:07 | 显示全部楼层
用户 发表于 2016-7-23 15:07
这是啪啪打google那个TPU的脸吧。之前google说一个order of magnitude的能耗比提升。老黄pascal x2倍int8 x ...

一看就是技术达人,请问“binary neuron”是什么意思?
发表于 2016-7-23 22:18 | 显示全部楼层
6666666666666cs本科毕业就工作了。看来应该去读读博士。可以名正言顺买titan了
发表于 2016-7-23 23:43 | 显示全部楼层
LZ显然需要召唤宽带树女神和D女王啊!
发表于 2016-7-24 05:14 | 显示全部楼层
本帖最后由 FancyMouse 于 2016-7-23 13:15 编辑

估摸着也就是一个乘法+聚集求和的指令,不用多少电路,强行op乘2。实际还得拿到新文档以后去研究。
不过如果真是这样,怎么不把and+popcnt算成给int1矩阵乘法加速的指令呢,还能强行op乘个8(
您需要登录后才可以回帖 登录 | 加入我们

本版积分规则

小黑屋|手机版|Archiver|Chiphell ( 沪ICP备12027953号-5 )沪公网备310112100042806

GMT+8, 2020-1-23 07:44 , Processed in 0.018165 second(s), 16 queries , Gzip On, Redis On.

Powered by Discuz! X3.1

© 2007-2019 Chiphell.com All rights reserved.

快速回复 返回顶部 返回列表