NVIDIA Tensor Core深度学习核心解析及跑分测试
核心解析
不久前，nvidia在siggraph 2018上正式发布了新一代gpu架构——turing（图灵），黄仁勋称turing架构是自2006年cuda gpu发明以来最大的飞跃。
turing架构的两大重要特性便是集成了用于光线追踪的rt core以及用于ai计算的tensor core，使其成为了全球首款支持实时光线追踪的gpu。
不过说到ai计算，nvidia gpu成为最好的加速器早已是公认的事实，但将tensor core印上gpu名片的并不是这次的turing，而是他的上任前辈——volta。
基于volta架构的titan v是nvidia在计算领域成就的集大成者。深度学习和神经网络已成为nvidia gpu的背后驱动力，作为最先进的计算加速器，它集成了用于机器学习操作的内置硬件和软件加速，深度学习能力完全可以被当做titan v和volta的名片。
titan v与初代基于开普勒的geforce gtx titan已经相去甚远，初代titan的定位是一款万能显卡，既可作为游戏发烧友的旗舰游戏显卡，也为专业消费者提供全双精度浮点（fp64）计算能力。
在titan v诞生之前，titan产品线几乎都是基于这种设计方法，一颗巨大的gpu核心是nvidia“高大全”设计思路的最好代表。
而在titan v上，nvidia再次扩展了大核心的上限。volta最引人注目的则是其全新的专用处理模块——tensor core（张量计算核心），它与volta的其他微架构改进，以及支持深度学习和高性能计算（hpc）的软件/框架集成在一起。
凭借面积达815mm?的巨大gv100核心，titan这一产品线变得比以往任何时候都更接近工作站级，titan v在拥有世界最强图形渲染性能的同时，深度学习和高性能计算方面的性能都有了极大的提升，当然它的价格也达到了工作站级的3000美元。
强化算力，重造titan
除了深度学习和通用计算之外，titan品牌的这次迭代还涉及到其他一些因素。 nvidia其实已经不太需要通过titan系列为自己树立形象，最初的gtx titan已经通过nvidia k20xs为oak ridge国家实验室的titan超级计算机提供计算力。况且，titan系列在产品价格和性能方面也没有任何特别的竞争压力。
尽管titan v的非ecc hbm2显存和geforce驱动程序堆栈都更加面向消费者，但该卡仍可直接受益于框架和api的软件支持，这是nvidia深度学习开发整体工作的一部分。
鉴于单路titan v并不会对服务器计算卡quadro gv100产生什么影响，nvidia在titan v上只砍掉了针对服务器多路互联设计的nvlink高速总线，而主要的计算能力（fp64/fp16/tensor core）都被完整保留。
与pascal及更早期的产品均使用面向pc设计的gpu以及常规gddr5(x)显存不同，nvidia这次选择了一颗规模巨大、产量和良品率都不高的服务器级芯片，有大量的晶体管被花费在了非图形功能上（即tensor core），这是nvidia在计算领域押下的赌注，nvidia已经不满足于只在传统图形计算卡和通用计算方面处于领先地位。
由于是首次分析gpu的深度学习性能，所以目前市面上还没有确定一套标准的基准测试，特别是对于volta独特的张量内核和混合精度功能。对于titan v，我们将使用百度deepbench、nvidia的caffe2 docker、stanford dawnbench和hpe深度学习基准套件（dlbs）来测试。
但是在深入研究这些测试数据之前，雷锋网首先会就深度学习、gpu、volta微架构以及深度学习性能基准进行一些背景介绍。
gpu与深度学习
首先要说明的是，虽然“机器学习”或更通用的“ai”有时可互换用于“深度学习”，但从技术上讲，它们各自指的是不同的东西，机器学习是ai的子集，深度学习则是机器学习的子集。
深度学习是因“深度神经网络”（deep neural networks）而得名，其最终被设计为识别数据中的模式，产生相关预测，接收关于预测准确度的反馈，然后基于反馈进行自我调整。计算发生在“节点”上，“节点”被组织成“层”：原始输入数据首先由“输入层”处理，“输出层”推出代表模型预测的数据。两者之间的任何一层都被称为“隐藏层”，而“deep”则代表着深度神经网络有许多隐藏层。
这些隐藏层可以在不断增加的抽象层次上运行，使得它们甚至可以从复杂的输入数据中提取和区分非线性特征。一个标准的例子是图像识别，其中初始层寻找某些边缘或形状，这通知后面的层寻找鼻子和眼睛，之后的层可能寻找面部。最后的图层组合了所有这些数据以进行分类。
随着输入数据在模型中向前推进，计算包括特殊的内部参数（权重），最后会产生一个表示模型预测与正确值之间误差的损失函数。然后使用此错误信息反向运行模型以计算将改善模型预测的权重调整，该前向和后向传递（或反向传播）序列包括单个训练迭代。
为了进行推断，这个过程自然地排除了逆向传递，最终需要的计算强度比训练模型更小。从这个意义上说，推断也不太需要像fp32这样高的精度，并且可以对模型进行适当的修剪和优化，以便在特定的设备上部署。然而推断设备对延迟、成本和功耗变得更加敏感，尤其是在边缘计算的场景下。
卷积神经网络（cnn）和递归神经网络（rnn）是深度神经网络的两个重要子类型。卷积本身是一种操作，将输入数据和卷积核结合起来形成某种特征映射，转换或过滤原始数据以提取特征。
cnn通常是“前馈”的，因为数据在没有循环的情况下流过各层。而对于rnn（以及像lstm和gru这样的变体）来讲，每次计算后都会有一个单独的权重循环回自身，给网络一种“记忆”感，这让网络能够做出有时间意识的预测，在文本分析等场景中很有用。
由于深度学习数学可以归结为线性代数，因此某些操作可以重写为对gpu更友好的矩阵间乘法。当nvidia首次开发并公布cudnn时，其中一个重要实现就是将算法降级为矩阵乘法以加速卷积。多年来cudnn的发展包括“预先计算的隐式gemm”卷积算法，它恰好是触发tensor core卷积加速的唯一算法。
nvidia gpu的优势
对于深度学习训练来说，gpu已经成为加速器的最佳选择。大多数计算本质上是并行的浮点计算，即大量的矩阵乘法，其最佳性能需要大量的内存带宽和大小，这些需求与hpc的需求非常一致，gpu正好可以提供高精度浮点计算、大量vram和并行计算能力，nvidia的cuda可谓恰逢其时。
cuda和nvidia的计算业务的发展与机器学习的研究进展相吻合，机器学习在2006年左右才重新成为“深度学习”。gpu加速神经网络模型相比cpu可提供数量级的加速，反过来又将深度学习重新推广到如今的流行词汇。与此同时，nvidia的图形竞争对手ati在2006年被amd收购；opencl 1.0在2009年才发布，同年amd剥离了他们的globalfoundries晶圆厂。
随着dl的研究人员和学者们成功地使用cuda来更快地训练神经网络模型，nvidia才发布了他们的cudnn库的优化深度学习原语，其中有很多以hpc为中心的blas（基本线性代数子例程）和相应的cublas先例，cudnn将研究人员创建和优化cuda代码以提高dl性能的需求抽象出来。至于amd的同类产品miopen，去年才在rocm保护伞下发布，目前也只在caffe公开发布。
所以从这个意义上讲，尽管nvidia和amd的底层硬件都适合dl加速，但nvidia gpu最终成为了深度学习的参考实现。
剖析tensor core
在关于volta混合精度tensor core的几个谜团中，一个比较烦人的问题是4 x 4矩阵乘法的能力。tensor core是一种新型处理核心，它执行一种专门的矩阵数学运算，适用于深度学习和某些类型的hpc。tensor core执行融合乘法加法，其中两个4*4 fp16矩阵相乘，然后将结果添加到4*4 fp16或fp32矩阵中，最终输出新的4*4 fp16或fp32矩阵。
nvidia将tensor core进行的这种运算称为混合精度数学，因为输入矩阵的精度为半精度，但乘积可以达到完全精度。碰巧的是，tensor core所做的这种运算在深度学习训练和推理中很常见。
tensor core虽然在gpu里是全新的运算单元，但其实它与标准的alu流水线并没有太大差别，只不过tensor core处理的是大型矩阵运算，而不是简单地单指令流多数据流标量运算。tensor core是灵活性和吞吐量权衡的选择，它在执行标量运算时的表现很糟糕，但它可以将更多的操作打包到同一个芯片区域。
tensor core虽然有一定的可编程性，但仍然停留在4*4矩阵乘法累加层面上，并且不清楚累积步骤是如何以及何时发生的。尽管被描述为进行4*4矩阵数学运算，但实际上tensor core运算似乎总是使用16*16矩阵，并且操作一次跨两个tensor core进行处理。这似乎与volta架构中的其他变化有关，更具体地说，与这些tensor core是如何集成进sm中有关。
对于volta架构，sm被划分为四个处理块或子核。对于每个子核，调度器每个时钟向本地分支单元（bru）、tensor core阵列、数学分派单元或共享mio单元发出一个warp指令，这就首先阻止了tensor运算和其他数学运算同时进行。在利用两个tensor core时，warp调度器直接发出矩阵乘法运算，并且在从寄存器接收输入矩阵之后，执行4*4*4矩阵乘法。待完成矩阵乘法后，tensor core再将得到的矩阵写回寄存器。
在tensor core执行实际指令时，即使在使用nvvm ir（llvm）的编译器级别上，也仅存在用于warp级矩阵操作的本征，对于cuda++和ptx isa，warp级别仍然是唯一级别。加载输入矩阵的形式是每个扭曲线程持有一个片段，其分布和身份均未指定。从广义上讲，它遵循标准cuda核心的基于线程级别拼接的gemm计算的相同模式。
一般而言，给定a*b+c tensor core操作，片段由a的8个fp16*2元素（即16个fp16元素）和b的另外8个fp16*2元素，以及fp16累加器的4个fp16*2元素或 fp32累加器的8个fp32元素组成。
在矩阵乘法累加运算之后，计算结果会分散在每个线程的目标寄存器片段中，需要在整个范围内统一，如果其中一个warp线程退出，这些低级操作基本上就会失败。
citadel llc团队的低级微基准测试揭示了许多volta微体系结构细节，包括tensor core操作和相关的片段，与输入矩阵相比，它们都位于寄存器和标识中。他们观察到，子核核心以特定的拼接模式计算矩阵乘法，其中所有32个warp线程都在运行。
从概念上讲，tensor core在4*4子矩阵上运行，以计算更大的16*16矩阵。warp线程被分成8组，每组4个线程，每个线程组连续计算一个8*4块，总共要经过4组的过程，每一个线程组都处理了目标矩阵的1/8。
在一个集合中，可以并行完成四个hmma步骤，每个步骤适用于4*2子块。这四个线程直接链接到寄存器中的那些矩阵值，因此线程组可以处理单个step 0 hmma指令，从而一次性计算子块。
由于矩阵乘法在数学上需要对某些行列进行复用，以允许跨所有8*4块并行执行，每个4*4矩阵被映射到两个线程的寄存器。在计算16*16父矩阵的4*4次子矩阵运算中，这将包括将连续计算的集合相加，形成16*16矩阵中4*8个元素的相应块。尽管citadel没有对fp16进行测试，但它们发现fp16 hmma指令只产生2个步骤，而不是4个步骤，这或许与fp16只占用的较小的寄存器空间有关。
通过独立的线程调度和执行，以及warp同步和warp-wide结果分配，基本的4*4*4 tensor core操作转换为半可编程16*16*16混合精度矩阵乘法累加。虽然cuda 9.1支持32*8*16 and 8*32*16矩阵，但相乘的矩阵都需要相应的列和行为16，最终矩阵为32*8或8*32。
tensor core的运行方式似乎是nvidia gemm计算层次结构的一个硬件实现的步骤，如cutlass（用于gemm操作的cuda c ++模板库）中所示。对于传统的cuda核心，最后一步需要将warp tile结构分解为由各个线程拥有的标量和向量元素。使用wmma api（现在表示张量核），所有这些都被抽象掉了，只剩下了需要处理的合作矩阵片段加载/存储和多重积累。积累发生在一个fma类型的操作中。
在寄存器级别上，nvidia在他们的hot chips 2017论文中提到“使用三个相对较小的乘法和累加器数据的4*4矩阵，可以执行64次乘加运算。”而增强的volta simt模型的每线程程序计数器（能够支持张量核）通常需要每个线程2个寄存器槽。hmma指令本身会尽可能多复用寄存器，所以我无法想象寄存器在大多数情况下不会出现瓶颈。
对于独立的4*4矩阵乘法累加，tensor core阵列在寄存器、数据路径和调度方面很有核能并没有物理设计，它只能用于特定的子矩阵乘法。
无论如何，从nvidia的角度来看，volta不是一颗深度学习的专用asic，它仍然覆盖gpgpu的领域，因此保持cuda可编程tensor core适用于gemm / cublas和hpc是合乎逻辑的。对于cuda c++的cutlass来说，情况更是如此，因为它的wmma api支持旨在为广泛的应用程序启用tensor coregemm操作。从根本上说，nvidia深度学习硬件加速的发展与cudnn（以及cublas）的发展有很大关系。
让fp16适用于深度学习
volta的深度学习能力是建立在利用半精度浮点（ieee-754 fp16）而非单精度浮点（fp32）进行深度学习训练的基础之上。
该能力首先由cudnn 3支持并在tegra x1的maxwell架构中实现，随后原生半精度计算被引入pascal架构并被称为“伪fp16”，即使用fp32 alus处理成对的fp16指令，理论上可以使每个时钟的fp16吞吐量增加一倍。这一特性实际上已经在tensor core处理寄存器中矩阵片段的过程中得到体现，其两个fp16输入矩阵被收集在8个fp16*2或16个fp16元素中。
就fp32与fp16而言，由于单精度浮点所包含的数据多于半精度浮点，因此计算量更大，需要更多的内存容量和带宽来容纳和传输数据，并带来更大的功耗。因此，在计算中成功使用低精度数据一直是穷人的圣杯，而目标则是那些不需要高精度数据的应用程序。
除了api/编译器/框架的支持之外，深度学习一直都有在使用fp16数据类型时损失精度的问题，这会让训练过程不够准确，模型无法收敛。
据雷锋网了解，nvidia以前也曾在类似的情况下推出过“混合精度”这一概念，在pascal的快速fp16（针对gp100）和dp4a/dp2a的整数点积操作（针对gp102、gp104和gp106 gpu）中，就曾提出过类似的说法。
当时人们关注的是推理能力，就像titan v的“深度学习tflops”一样，titan x（pascal）推出了“44 tops（新型深度学习推断指令）”。新的指令对4元8位向量或2元8位/16位向量执行整数点积，从而得到一个32位整数积，可以与其他32位整数一起累积。
对于volta的混合精度而言，重要的精度敏感数据（如主权重）仍然会存储为fp32；而tensor core的混合精度训练则会将两个半精度输入矩阵相乘得到一个精度乘积，然后累积成一个精度和。nvidia表示，在将结果写入内存之前，tensor core会将结果转换回半精度，这样当使用半精度格式时，寄存器和存储器中需要的数据更少，这有助于处理超大矩阵中的数据。
fp16与fp32所包含的数据空间并不相同，归一化方法可以解决fp32格式数据超出fp16可表示范围的问题。举个例子，许多激活梯度的值都落在fp16的范围之外，但由于这些值聚集在一起，因此将损耗乘以缩放因子便可以移动fp16范围内的大部分值。在完成最终的权重更新之前，将梯度重新缩放到原始范围，便可以维持其原始的精度。
不过，并非所有的数学、神经网络和层都适用于fp16，通常fp16和tensor core的混合精度最适合卷积和rnn重图像处理等，而对于不适合的神经网络框架或类型，fp16将默认禁用或不推荐使用。
内存改进，sm变化
使用tensor core处理混合精度数据似乎可以减轻内存带宽问题，但事实证明，尽管volta在几乎所有方面都得到了内存子系统的增强，但幅度并不明显。
首先，volta有一个12 kb的l0指令缓存，虽然pascal及其他之前的gpu核心也有指令缓存，但volta更高效的l0是子核心sm分区私有的，因此它对warp调度器来说也是私有的，这对volta架构更大的指令大小是一种补偿，并有可能为支持tensor core通道的框架做出贡献。同时volta指令延迟也要低于pascal，特别是核心fmas从6个周期减少到了4个周期。
随着每个sm中调度器比例的增加，砍掉第二个调度端口似乎是对具有独立数据路径和数学调度单元的子核心的权衡。而具备fp32/int32执行能力，也为其他低精度/混合精度模型打开了大门。这些子核方面的增强，都是为了优化tensor core阵列。
另一个重大变化是合并l1缓存和共享内存。在同一个块中，共享内存可配置为每sm 最高96 kb。hbm2控制器也进行了更新，其效率提高了10~15%。
深度学习基准测试
深度学习从框架到模型，再到api和库，ai硬件的许多部分都是高度定制化的，这样的新领域有时会让人非常难以理解。
俗话说“光说不练假把式”，实践永远是检验真理的唯一标准。对计算机来说，介绍的再详细也不如真刀真枪跑一下测试，没有什么比benchmark更能提现硬件的实际表现了。
随着imagenet和一些衍生模型（alexnet、vggnet、inception、resnet等）的影响，ilsvrc2012（imagenet大规模视觉识别挑战）中的图像数据集训练逐渐被行业所认可。
现在基本上所有深度学习框架都支持cuda和cudnn，对于volta而言，支持fp16存储的框架也都支持tensor core加速，启用fp16存储后tensor core加速会自动启用。
在随后的下篇文章里，我们将通过多项测试来考验volta架构，利用各种深度学习框架来了解tensor core的性能。
测试篇
本篇将通过多项测试来考验volta架构，利用各种深度学习框架来了解tensor core的性能。
深度学习的基准测试
很多时候，深度学习这样的新领域会让人难以理解。从框架到模型，再到api和库，ai硬件的许多部分都是高度定制化的，因而被行业接受的公开基准测试工具很少也就不足为奇。
随着imagenet和一些衍生模型（alexnet、vggnet、inception、resnet等）的影响，ilsvrc2012（imagenet大规模视觉识别挑战）中的图像数据集训练逐渐被行业所认可。
基本上所有现代深度学习框架都支持cuda和cudnn，对于volta而言，所有支持fp16存储的框架也都支持tensor core加速，启用fp16存储后tensor core加速会自动启用，因此我们可以利用这些框架来了解tensor core的性能。
在常见的第三方深度学习基准套件中，fathom和tbd是更传统的基准测试套件，其测试针对特定框架和模型进行了配置，涵盖了许多不同的机器学习应用程序。 同时，最近的深度学习框架侧重于比较给定模型和跨框架的数据集的性能。
而deepbench本身并不使用框架，而是使用低级库来评估不同设备的机器学习性能。就其本身而言，虽然它并不直接将框架/模型/应用程序性能与其他测试联系在一起，但它提供了代表供应商优化的数学操作和硬件性能的指标，每个产品的二进制文件都使用硬件供应商提供的库进行编译。
dawnbench则更加与众不同，与其说它是一个基准测试套件，不如说是对三个数据集（imagenet、cifar10和squad）的训练和推断结果进行类似于竞赛的报告，重点考量端对端的计算精确度和成本。
至于hpe dlbs，作为hpe深度学习指南的一部分，它主要以gpu为中心，坚持使用tensorflow、mxnet、pytorch和caffe类型框架，还包括tensorrt测试。虽然其具有良好的多测试批处理、日志记录、监控和报告功能，但它只输出纯粹的性能和时间指标，不涉及端对端的时间精度或成本。
从这些基准测试软件中可以看出，深度学习框架之间的差异很容易使测试结果变得毫无意义，从而影响我们对这些框架的研究。convnet-benchmark和pytorch的创始人soumith chintala指出，如果没有机器学习的背景，很难独立地验证深度学习基准测试的准确性和范围，不过mlperf测试项目似乎试图解决这个问题。
mlperf是由dawnbench等测试软件的设计者和工程师联合打造的全新高端基准测试套件，希望囊括fathom的跨域测试方法以及dawnbench对超过阈值精度模型的端对端计算时间考察。不过它目前正在处于alpha阶段，开发团队表示其尚不适合进行精确的硬件对比。
综合考虑之下，本次测试将不包含mlperf项目，而是使用deepbench、caffe2 docke、stanford dawn和hpe dlbs来进行。
deepbench训练测试之gemm和rnn
首先进行的是gemm测试，利用某些深度学习应用程序（deepspeech、speaker id和language modeling）中的内核进行gemm操作，测出的性能比在cublas中运行纯矩阵-矩阵乘法更有代表性。
测试的结果在意料之内，启用tensor core可以大幅提升性能。深入研究细节可以发现，tensor core对于特定类型的矩阵-矩阵乘法会有特别的影响。
通过深度学习应用程序拆分gemm测试，我们可以了解tensor core在理想和非理想情况下的表现。
speaker id gemm工作负载实际上只包含两个内核，其中10微秒的时间差意味着大约1 tflops的算力差异。
通过对语言模型内核的研究，可以了解tensor core在非理想情况下的性能。这些核矩阵的大小是m=512或1024，n=8或16，k=500000，虽然每个数在技术上都可以被8整除——这是满足张量核加速度的基本要求之一——但这些矩阵的形状与tensor core支持的16*16*16、32*8*16和8*32*16等基本wmma形状不太匹配。假如tensor core真正在独立的8x8x8级别上运行，那么运算8*8*8矩阵的性能也不会很好。
因此，tensor core无法高效的将这些非常不平衡的矩阵分解为n=8或16。而且，tensor core在deepspeech内核上的性能也出现异常：
从所有子项的平均成绩来看，这个浮点运算性能令人印象深刻。当矩阵适合于tensor core时，性能可以超过90tflops；相反如果二者无法契合，并正确的换位没有发挥作用，性能会低至<1tflops的水平。
对于deepbench rnn内核的测试，rnn类型之间没有明显的差异，但是在每种rnn类型中，如果将不同内核挨个进行对比判断，也可以看到与gemm中相同的趋势。
比较有趣的是，titan xp与titan v在未使用tensor core加速时的表现有很接近，titan xp的高频率为其性能起到了一定的帮助。
deepbench训练测试之convolutions
在卷积训练工作负载测试中，tensor core再次显着提高了性能。鉴于卷积层是图像识别和分类的基础，因而卷积运算是tensor core加速的最大潜在受益者之一。
从所有测试项的平均成绩可以看出，volta在启用了tensor core的fp16混合精度运算能力后性能再次取得了领先。不过与gemm不同，在fp32卷积上启用tensor core会导致明显的性能损失。
当计算涉及不匹配的张量尺寸时，标准精度模式遵循cudnn指定的最快前向算法（如winograd），而混合精度模式必须对所有内核使用隐式预计算gemm，这会造成两种混合精度模式的性能会出现下滑。
要符合tensor core加速的要求，输入和输出通道尺寸必须是8的倍数，输入、过滤和输出数据的类型必须是半精度。使用tensor core实现卷积加速要求张量采用nhwc格式，但大多数框架都希望采用nchw格式的张量。在这种情况下，输入通道不是8的倍数，但测试程序会自动填充以解决此问题。
需要注意的是，所有这些nchw内核都需要转换为nhwc。想要从tensor core中受益，需要正确的调整卷积格式，本次测试使用的是nvidia提供的标准库和makefile。nvidia指出，一旦进行加速卷积，它会消耗掉相当多的运行时间，这将会对fp32和fp16混合精度模式造成影响。
deepbench推理测试之gemm
数据精度方面，百度将deepbench gemm和卷积定义支持32bit累加的int8格式，以支持volta和pascal上的int8计算。
titan v和titan xp均拥有4倍于int32的int8性能，deepbench的int8推理测试正中pascal引入的dp4a矢量点积能力之下怀。volta同样拥有这一能力，在指令集中二者均显示为idp和idp4a。
对igemm来说，正如cutlass所示，dp4a是一项定制操作。因此除语言建模之外，int8的性能都非常之高。当然，与硬件不匹配的张量尺寸不适合tensor core加速，这一点与之前完全一样。
在完全连接（仿射）层中，每个节点都与前一层中的各节点相连接。对于一个典型的cnn来说，完全连接的层意味着可以结合所有提取的特征做出最终预测并对图像进行分类。这些测试结果数据也意味着大型且规则的矩阵可以在tensor core加速中获得更大的收益。
deepbench推理测试之convolutions
再次来到卷积测试环节，8位乘法/32位累加再次出现在int8推理中。
测试中最引人注目的是titan xp，在resnet、speaker id和vision项目中，titan xp表现出了强劲的int8吞吐量。
从内核方面来看，并没有发现这一现象的根源所在，猜测可能是由于pascal的dp4a库好驱动程序比volta更为成熟所致，亦或许是volta通过单独的int单元处理这些运算。
deepbench推理测试之rnn和sparse gemm
deepbench的最后一项推理测试是rnn和sparse gemm，虽然测试中可以选择fp16，但实际上它们都只支持fp32运算。
虽然rnn可能会有加速，但deepbench和nvidia目前仅支持单精度rnn推理。
nvidia caffe2测试之resnet50和imagenet
虽然内核和深度学习数学运算可能很有用，但实际应用中是使用真实数据集进行训练的。使用标准的ilsvrc 2012图片集，在imagenet上通过resnet50模型来训练和推断，可以展示更具参考的性能数据。
虽然fp16和tensor core有单独的开关，但titan v在启用和禁用tensor core的情况下运行fp16的性能是完全一样的。
只看原始吞吐量性能的话，titan v在所有批尺寸下都处于领先地位。凭借tensor core，titan v可处理的批尺寸达到甚至超过了64，而其他显卡即便有12 gb显存也无法与之相比。
不过只看原始吞吐量性能的问题在于，深度学习的实际性能从来没有这么简单。首先，许多模型可能会牺牲精度和训练时间以换取针对吞吐量的优化，如果模型需要较长的时间来收敛，那么每秒训练的峰值性能就没有参考意义了。
这些问题与使用fp16存储和tensor core的volta尤为相关，如果在实际应用中使用了有损缩放或单精度批量归一化，这在吞吐量性能中都是无法体现的。
hpe dlbs caffe2测试之resnet50和imagenet
接下来，我们看一下深度学习指南中的hpe dlbs。与通常的深度学习测试不同，hpe dlbs基本上只输出吞吐量和时间指标。
hpe dlbs的一大特色是支持nvidia caffe2测试使用的数据集，我们同样可以在imagenet上使用resnet50模型来训练和推断。但是由于二者的模型和实现不同，测出的吞吐量性能无法与nvidia caffe2直接进行比较。
在测试中，titan v无法支持某些特定的批尺寸，但总体趋势和之前的测试基本相同，fp16和tensor core提供了更高的吞吐量。不过遗憾的是，hpe dlbs caffe2测试似乎不支持int8推理。
hpe dlbs tensorrt测试之resnet50和imagenet
hpe dlbs的另一大特色是支持tensorrt（nvidia推理优化引擎）的基准测试功能， nvidia近年来已将tensorrt与新的深度学习功能（如int8/dp4a和tensor core的16位累加器模式）相结合以进行推理。
使用caffe模型，tensorrt可以根据需要调整模型，以便在给定的精度下进行推理。我们在titan x（maxwell）和titan xp（pascal）上运行了64、512和1024的批尺寸，在titan v运行了128、256和640的批尺寸。
titan xp的高int8性能在一定程度上印证了gemm/卷积性能，这两个工作负载似乎都在使用dp4a。不过雷锋网(公众号：雷锋网)并未了解到dp4a如何在titan v上实现，只知道它由volta指令集提供支持，且volta确实拥有独立的int32单元。
dawnbench测试之cifar10图像分类
就实际应用的性能而言，深度学习训练更适合用时间/准确性和成本来描述，而这两点分别对应dawnbench的两项子测试。对于使用cifar10的图像分类来说，这两项子测试为：
时间/准确性：训练cifar10数据集的图像分类模型，报告训练所需的时间，且要求测试集准确性至少为94%。
成本：在公共云基础架构上，计算达到94％或更高的测试集准确性所需的总时间，将所花费的时间（以小时为单位）乘以每小时实例的成本，以获得训练模型的总成本。
测试选用pytorch的cifar10训练实现中最快的两个，其中一个基于resnet34，是为了在nvidia geforce gtx 1080 ti上运行，而第二个基于resnet18，是为了在单个tesla v100上运行。这些都是dawnbench中最近的热门测试，可以认为它们是相当符合现代的项目，同时cifar10也不是一个非常密集的数据集。
cifar10的小型图像数据集运行良好，第一个训练实现是在单个gtx 1080 ti上运行，需要35分37秒才能训练到94％的准确性，而在第二个训练实现中，titan v只用了5分41秒就完成了94%的目标。
顺带一提，虽然titan v在第一个训练实现中不会使用tensor core，但凭借相对于pascal的一般改进，titan v在这个测试中的速度依然比titan xp快20％左右，同时系统峰值功耗也下降了大约80w。
结语
tensor core是titan v的重要组成部分，本文的目的也是尽可能的了解volta，所以测试着重考察了tensor core加速。
本次测试还有许多没有涉及到的各种其他测试和套件，它们普遍在设计上有太多欠缺。事实证明，无论从整体还是局部来看，都不可能找到一个负载贴合当下实际、提供端到端指标、覆盖多个机器学习域、支持tensor core和混合精度，最重要的是易于非专业开发者使用的深度学习套件。
即便是参考价值较大的dawnbench，设计本意也并非提供一个通用的基准，而是为方便研究者和开发者创建他们自己的实现。dawnbench的深度学习框架仍然需要修改以作为一款有效的基准测试使用，但重新配置一个与volta兼容的混合精度模型就不是一件可以轻松完成的事情。
这实际上与titan v自身相关，tensor core和混合精度需要专门的开发来适配支持，只能在特定情况下明显提升性能，而且即便用户的代码和数据集可以与tensor core良好的匹配，到头来神经网络处理也会受到来自传统alu的限制。
而对于主流消费者群体来说，tensor core的发展对他们意味着什么？
最新的turing架构证明，tensor core在游戏卡领域同样可以有所应用。rtx光线追踪技术就是使用tensor core对图像进行降噪以弥补光线数量的限制。nvidia唯一需要考虑的问题是，这些玩家是否乐意为这些tensor core买单。
对于任何考虑购买titan v以满足计算需求的人来说，投资titan v就意味着投资混合精度深度学习模型和基于wmma的hpc gemm加速。在cudnn、cublas以及早期dp4a和fp16*2混合精度计算的背景下，tensor core是试图用可编程硬件来实现深度学习加速的一种自然进化。
可以肯定的是，titan v绝对代表了nvidia对未来gpgpu的愿望。
2021中国汽车供应链峰会暨第六届铃轩奖盛典圆满落幕
数据网格是什么意思？为什么要考虑实现数据网格？
Anritsu MS8604A数字移动无线发射机测试仪孟姐/
结合从业者Michail Katkoff的数据分享以及《一起来捉妖》的游戏特色，来分析该品类的现状和发展趋势
GCAN以太网can网关如何修改本地IP？
NVIDIA Tensor Core深度学习核心解析及跑分测试
debian软件包管理详解
ADI公司推出集成式隔离电源控制器系列
国民经济的基础性行业之一---减速机行业
亚马逊涉嫌未能保护工人免受COVID-19的侵害而被起诉
无位置反馈比例阀放大器RT-PVDA
AMD官方澄清RX5700XT50周年纪念版显卡限地区发售的情况
全球领先系统制造商推出 NVIDIA AI 就绪型服务器，为企业大幅提升生成式 AI 性能
ds18b20如何读取温度？温度数据怎么换算？
Linux操作系统的系统性能监控工具syastat的详解
首创后置三摄+三摄万物：华为联合京东这波海报666
物联网安全之Unicorn框架初探
高频闭锁方向保护的动作条件和闭锁条件是什么？
七月新机红米Pro2，屏幕占比史无前例，价格适中1999
华为Mate50“没电也能打电话”
不久前，nvidia在siggraph 2018上正式发布了新一代gpu架构——turing（图灵），黄仁勋称turing架构是自2006年cuda gpu发明以来最大的飞跃。
turing架构的两大重要特性便是集成了用于光线追踪的rt core以及用于ai计算的tensor core，使其成为了全球首款支持实时光线追踪的gpu。
不过说到ai计算，nvidia gpu成为最好的加速器早已是公认的事实，但将tensor core印上gpu名片的并不是这次的turing，而是他的上任前辈——volta。
基于volta架构的titan v是nvidia在计算领域成就的集大成者。深度学习和神经网络已成为nvidia gpu的背后驱动力，作为最先进的计算加速器，它集成了用于机器学习操作的内置硬件和软件加速，深度学习能力完全可以被当做titan v和volta的名片。
titan v与初代基于开普勒的geforce gtx titan已经相去甚远，初代titan的定位是一款万能显卡，既可作为游戏发烧友的旗舰游戏显卡，也为专业消费者提供全双精度浮点（fp64）计算能力。
在titan v诞生之前，titan产品线几乎都是基于这种设计方法，一颗巨大的gpu核心是nvidia“高大全”设计思路的最好代表。
而在titan v上，nvidia再次扩展了大核心的上限。volta最引人注目的则是其全新的专用处理模块——tensor core（张量计算核心），它与volta的其他微架构改进，以及支持深度学习和高性能计算（hpc）的软件/框架集成在一起。
凭借面积达815mm?的巨大gv100核心，titan这一产品线变得比以往任何时候都更接近工作站级，titan v在拥有世界最强图形渲染性能的同时，深度学习和高性能计算方面的性能都有了极大的提升，当然它的价格也达到了工作站级的3000美元。
强化算力，重造titan
除了深度学习和通用计算之外，titan品牌的这次迭代还涉及到其他一些因素。 nvidia其实已经不太需要通过titan系列为自己树立形象，最初的gtx titan已经通过nvidia k20xs为oak ridge国家实验室的titan超级计算机提供计算力。况且，titan系列在产品价格和性能方面也没有任何特别的竞争压力。
尽管titan v的非ecc hbm2显存和geforce驱动程序堆栈都更加面向消费者，但该卡仍可直接受益于框架和api的软件支持，这是nvidia深度学习开发整体工作的一部分。
鉴于单路titan v并不会对服务器计算卡quadro gv100产生什么影响，nvidia在titan v上只砍掉了针对服务器多路互联设计的nvlink高速总线，而主要的计算能力（fp64/fp16/tensor core）都被完整保留。
与pascal及更早期的产品均使用面向pc设计的gpu以及常规gddr5(x)显存不同，nvidia这次选择了一颗规模巨大、产量和良品率都不高的服务器级芯片，有大量的晶体管被花费在了非图形功能上（即tensor core），这是nvidia在计算领域押下的赌注，nvidia已经不满足于只在传统图形计算卡和通用计算方面处于领先地位。
由于是首次分析gpu的深度学习性能，所以目前市面上还没有确定一套标准的基准测试，特别是对于volta独特的张量内核和混合精度功能。对于titan v，我们将使用百度deepbench、nvidia的caffe2 docker、stanford dawnbench和hpe深度学习基准套件（dlbs）来测试。
但是在深入研究这些测试数据之前，雷锋网首先会就深度学习、gpu、volta微架构以及深度学习性能基准进行一些背景介绍。
gpu与深度学习
首先要说明的是，虽然“机器学习”或更通用的“ai”有时可互换用于“深度学习”，但从技术上讲，它们各自指的是不同的东西，机器学习是ai的子集，深度学习则是机器学习的子集。
深度学习是因“深度神经网络”（deep neural networks）而得名，其最终被设计为识别数据中的模式，产生相关预测，接收关于预测准确度的反馈，然后基于反馈进行自我调整。计算发生在“节点”上，“节点”被组织成“层”：原始输入数据首先由“输入层”处理，“输出层”推出代表模型预测的数据。两者之间的任何一层都被称为“隐藏层”，而“deep”则代表着深度神经网络有许多隐藏层。
这些隐藏层可以在不断增加的抽象层次上运行，使得它们甚至可以从复杂的输入数据中提取和区分非线性特征。一个标准的例子是图像识别，其中初始层寻找某些边缘或形状，这通知后面的层寻找鼻子和眼睛，之后的层可能寻找面部。最后的图层组合了所有这些数据以进行分类。
随着输入数据在模型中向前推进，计算包括特殊的内部参数（权重），最后会产生一个表示模型预测与正确值之间误差的损失函数。然后使用此错误信息反向运行模型以计算将改善模型预测的权重调整，该前向和后向传递（或反向传播）序列包括单个训练迭代。
为了进行推断，这个过程自然地排除了逆向传递，最终需要的计算强度比训练模型更小。从这个意义上说，推断也不太需要像fp32这样高的精度，并且可以对模型进行适当的修剪和优化，以便在特定的设备上部署。然而推断设备对延迟、成本和功耗变得更加敏感，尤其是在边缘计算的场景下。
卷积神经网络（cnn）和递归神经网络（rnn）是深度神经网络的两个重要子类型。卷积本身是一种操作，将输入数据和卷积核结合起来形成某种特征映射，转换或过滤原始数据以提取特征。
cnn通常是“前馈”的，因为数据在没有循环的情况下流过各层。而对于rnn（以及像lstm和gru这样的变体）来讲，每次计算后都会有一个单独的权重循环回自身，给网络一种“记忆”感，这让网络能够做出有时间意识的预测，在文本分析等场景中很有用。
由于深度学习数学可以归结为线性代数，因此某些操作可以重写为对gpu更友好的矩阵间乘法。当nvidia首次开发并公布cudnn时，其中一个重要实现就是将算法降级为矩阵乘法以加速卷积。多年来cudnn的发展包括“预先计算的隐式gemm”卷积算法，它恰好是触发tensor core卷积加速的唯一算法。
nvidia gpu的优势
对于深度学习训练来说，gpu已经成为加速器的最佳选择。大多数计算本质上是并行的浮点计算，即大量的矩阵乘法，其最佳性能需要大量的内存带宽和大小，这些需求与hpc的需求非常一致，gpu正好可以提供高精度浮点计算、大量vram和并行计算能力，nvidia的cuda可谓恰逢其时。
cuda和nvidia的计算业务的发展与机器学习的研究进展相吻合，机器学习在2006年左右才重新成为“深度学习”。gpu加速神经网络模型相比cpu可提供数量级的加速，反过来又将深度学习重新推广到如今的流行词汇。与此同时，nvidia的图形竞争对手ati在2006年被amd收购；opencl 1.0在2009年才发布，同年amd剥离了他们的globalfoundries晶圆厂。
随着dl的研究人员和学者们成功地使用cuda来更快地训练神经网络模型，nvidia才发布了他们的cudnn库的优化深度学习原语，其中有很多以hpc为中心的blas（基本线性代数子例程）和相应的cublas先例，cudnn将研究人员创建和优化cuda代码以提高dl性能的需求抽象出来。至于amd的同类产品miopen，去年才在rocm保护伞下发布，目前也只在caffe公开发布。
所以从这个意义上讲，尽管nvidia和amd的底层硬件都适合dl加速，但nvidia gpu最终成为了深度学习的参考实现。
剖析tensor core
在关于volta混合精度tensor core的几个谜团中，一个比较烦人的问题是4 x 4矩阵乘法的能力。tensor core是一种新型处理核心，它执行一种专门的矩阵数学运算，适用于深度学习和某些类型的hpc。tensor core执行融合乘法加法，其中两个4*4 fp16矩阵相乘，然后将结果添加到4*4 fp16或fp32矩阵中，最终输出新的4*4 fp16或fp32矩阵。
nvidia将tensor core进行的这种运算称为混合精度数学，因为输入矩阵的精度为半精度，但乘积可以达到完全精度。碰巧的是，tensor core所做的这种运算在深度学习训练和推理中很常见。
tensor core虽然在gpu里是全新的运算单元，但其实它与标准的alu流水线并没有太大差别，只不过tensor core处理的是大型矩阵运算，而不是简单地单指令流多数据流标量运算。tensor core是灵活性和吞吐量权衡的选择，它在执行标量运算时的表现很糟糕，但它可以将更多的操作打包到同一个芯片区域。
tensor core虽然有一定的可编程性，但仍然停留在4*4矩阵乘法累加层面上，并且不清楚累积步骤是如何以及何时发生的。尽管被描述为进行4*4矩阵数学运算，但实际上tensor core运算似乎总是使用16*16矩阵，并且操作一次跨两个tensor core进行处理。这似乎与volta架构中的其他变化有关，更具体地说，与这些tensor core是如何集成进sm中有关。
对于volta架构，sm被划分为四个处理块或子核。对于每个子核，调度器每个时钟向本地分支单元（bru）、tensor core阵列、数学分派单元或共享mio单元发出一个warp指令，这就首先阻止了tensor运算和其他数学运算同时进行。在利用两个tensor core时，warp调度器直接发出矩阵乘法运算，并且在从寄存器接收输入矩阵之后，执行4*4*4矩阵乘法。待完成矩阵乘法后，tensor core再将得到的矩阵写回寄存器。
在tensor core执行实际指令时，即使在使用nvvm ir（llvm）的编译器级别上，也仅存在用于warp级矩阵操作的本征，对于cuda++和ptx isa，warp级别仍然是唯一级别。加载输入矩阵的形式是每个扭曲线程持有一个片段，其分布和身份均未指定。从广义上讲，它遵循标准cuda核心的基于线程级别拼接的gemm计算的相同模式。
一般而言，给定a*b+c tensor core操作，片段由a的8个fp16*2元素（即16个fp16元素）和b的另外8个fp16*2元素，以及fp16累加器的4个fp16*2元素或 fp32累加器的8个fp32元素组成。
在矩阵乘法累加运算之后，计算结果会分散在每个线程的目标寄存器片段中，需要在整个范围内统一，如果其中一个warp线程退出，这些低级操作基本上就会失败。
citadel llc团队的低级微基准测试揭示了许多volta微体系结构细节，包括tensor core操作和相关的片段，与输入矩阵相比，它们都位于寄存器和标识中。他们观察到，子核核心以特定的拼接模式计算矩阵乘法，其中所有32个warp线程都在运行。
从概念上讲，tensor core在4*4子矩阵上运行，以计算更大的16*16矩阵。warp线程被分成8组，每组4个线程，每个线程组连续计算一个8*4块，总共要经过4组的过程，每一个线程组都处理了目标矩阵的1/8。
在一个集合中，可以并行完成四个hmma步骤，每个步骤适用于4*2子块。这四个线程直接链接到寄存器中的那些矩阵值，因此线程组可以处理单个step 0 hmma指令，从而一次性计算子块。
由于矩阵乘法在数学上需要对某些行列进行复用，以允许跨所有8*4块并行执行，每个4*4矩阵被映射到两个线程的寄存器。在计算16*16父矩阵的4*4次子矩阵运算中，这将包括将连续计算的集合相加，形成16*16矩阵中4*8个元素的相应块。尽管citadel没有对fp16进行测试，但它们发现fp16 hmma指令只产生2个步骤，而不是4个步骤，这或许与fp16只占用的较小的寄存器空间有关。
通过独立的线程调度和执行，以及warp同步和warp-wide结果分配，基本的4*4*4 tensor core操作转换为半可编程16*16*16混合精度矩阵乘法累加。虽然cuda 9.1支持32*8*16 and 8*32*16矩阵，但相乘的矩阵都需要相应的列和行为16，最终矩阵为32*8或8*32。
tensor core的运行方式似乎是nvidia gemm计算层次结构的一个硬件实现的步骤，如cutlass（用于gemm操作的cuda c ++模板库）中所示。对于传统的cuda核心，最后一步需要将warp tile结构分解为由各个线程拥有的标量和向量元素。使用wmma api（现在表示张量核），所有这些都被抽象掉了，只剩下了需要处理的合作矩阵片段加载/存储和多重积累。积累发生在一个fma类型的操作中。
在寄存器级别上，nvidia在他们的hot chips 2017论文中提到“使用三个相对较小的乘法和累加器数据的4*4矩阵，可以执行64次乘加运算。”而增强的volta simt模型的每线程程序计数器（能够支持张量核）通常需要每个线程2个寄存器槽。hmma指令本身会尽可能多复用寄存器，所以我无法想象寄存器在大多数情况下不会出现瓶颈。
对于独立的4*4矩阵乘法累加，tensor core阵列在寄存器、数据路径和调度方面很有核能并没有物理设计，它只能用于特定的子矩阵乘法。
无论如何，从nvidia的角度来看，volta不是一颗深度学习的专用asic，它仍然覆盖gpgpu的领域，因此保持cuda可编程tensor core适用于gemm / cublas和hpc是合乎逻辑的。对于cuda c++的cutlass来说，情况更是如此，因为它的wmma api支持旨在为广泛的应用程序启用tensor coregemm操作。从根本上说，nvidia深度学习硬件加速的发展与cudnn（以及cublas）的发展有很大关系。
让fp16适用于深度学习
volta的深度学习能力是建立在利用半精度浮点（ieee-754 fp16）而非单精度浮点（fp32）进行深度学习训练的基础之上。
该能力首先由cudnn 3支持并在tegra x1的maxwell架构中实现，随后原生半精度计算被引入pascal架构并被称为“伪fp16”，即使用fp32 alus处理成对的fp16指令，理论上可以使每个时钟的fp16吞吐量增加一倍。这一特性实际上已经在tensor core处理寄存器中矩阵片段的过程中得到体现，其两个fp16输入矩阵被收集在8个fp16*2或16个fp16元素中。
就fp32与fp16而言，由于单精度浮点所包含的数据多于半精度浮点，因此计算量更大，需要更多的内存容量和带宽来容纳和传输数据，并带来更大的功耗。因此，在计算中成功使用低精度数据一直是穷人的圣杯，而目标则是那些不需要高精度数据的应用程序。
除了api/编译器/框架的支持之外，深度学习一直都有在使用fp16数据类型时损失精度的问题，这会让训练过程不够准确，模型无法收敛。
据雷锋网了解，nvidia以前也曾在类似的情况下推出过“混合精度”这一概念，在pascal的快速fp16（针对gp100）和dp4a/dp2a的整数点积操作（针对gp102、gp104和gp106 gpu）中，就曾提出过类似的说法。
当时人们关注的是推理能力，就像titan v的“深度学习tflops”一样，titan x（pascal）推出了“44 tops（新型深度学习推断指令）”。新的指令对4元8位向量或2元8位/16位向量执行整数点积，从而得到一个32位整数积，可以与其他32位整数一起累积。
对于volta的混合精度而言，重要的精度敏感数据（如主权重）仍然会存储为fp32；而tensor core的混合精度训练则会将两个半精度输入矩阵相乘得到一个精度乘积，然后累积成一个精度和。nvidia表示，在将结果写入内存之前，tensor core会将结果转换回半精度，这样当使用半精度格式时，寄存器和存储器中需要的数据更少，这有助于处理超大矩阵中的数据。
fp16与fp32所包含的数据空间并不相同，归一化方法可以解决fp32格式数据超出fp16可表示范围的问题。举个例子，许多激活梯度的值都落在fp16的范围之外，但由于这些值聚集在一起，因此将损耗乘以缩放因子便可以移动fp16范围内的大部分值。在完成最终的权重更新之前，将梯度重新缩放到原始范围，便可以维持其原始的精度。
不过，并非所有的数学、神经网络和层都适用于fp16，通常fp16和tensor core的混合精度最适合卷积和rnn重图像处理等，而对于不适合的神经网络框架或类型，fp16将默认禁用或不推荐使用。
内存改进，sm变化
使用tensor core处理混合精度数据似乎可以减轻内存带宽问题，但事实证明，尽管volta在几乎所有方面都得到了内存子系统的增强，但幅度并不明显。
首先，volta有一个12 kb的l0指令缓存，虽然pascal及其他之前的gpu核心也有指令缓存，但volta更高效的l0是子核心sm分区私有的，因此它对warp调度器来说也是私有的，这对volta架构更大的指令大小是一种补偿，并有可能为支持tensor core通道的框架做出贡献。同时volta指令延迟也要低于pascal，特别是核心fmas从6个周期减少到了4个周期。
随着每个sm中调度器比例的增加，砍掉第二个调度端口似乎是对具有独立数据路径和数学调度单元的子核心的权衡。而具备fp32/int32执行能力，也为其他低精度/混合精度模型打开了大门。这些子核方面的增强，都是为了优化tensor core阵列。
另一个重大变化是合并l1缓存和共享内存。在同一个块中，共享内存可配置为每sm 最高96 kb。hbm2控制器也进行了更新，其效率提高了10~15%。
深度学习基准测试
深度学习从框架到模型，再到api和库，ai硬件的许多部分都是高度定制化的，这样的新领域有时会让人非常难以理解。
俗话说“光说不练假把式”，实践永远是检验真理的唯一标准。对计算机来说，介绍的再详细也不如真刀真枪跑一下测试，没有什么比benchmark更能提现硬件的实际表现了。
随着imagenet和一些衍生模型（alexnet、vggnet、inception、resnet等）的影响，ilsvrc2012（imagenet大规模视觉识别挑战）中的图像数据集训练逐渐被行业所认可。
现在基本上所有深度学习框架都支持cuda和cudnn，对于volta而言，支持fp16存储的框架也都支持tensor core加速，启用fp16存储后tensor core加速会自动启用。
在随后的下篇文章里，我们将通过多项测试来考验volta架构，利用各种深度学习框架来了解tensor core的性能。
测试篇
本篇将通过多项测试来考验volta架构，利用各种深度学习框架来了解tensor core的性能。
深度学习的基准测试
很多时候，深度学习这样的新领域会让人难以理解。从框架到模型，再到api和库，ai硬件的许多部分都是高度定制化的，因而被行业接受的公开基准测试工具很少也就不足为奇。
随着imagenet和一些衍生模型（alexnet、vggnet、inception、resnet等）的影响，ilsvrc2012（imagenet大规模视觉识别挑战）中的图像数据集训练逐渐被行业所认可。
基本上所有现代深度学习框架都支持cuda和cudnn，对于volta而言，所有支持fp16存储的框架也都支持tensor core加速，启用fp16存储后tensor core加速会自动启用，因此我们可以利用这些框架来了解tensor core的性能。
在常见的第三方深度学习基准套件中，fathom和tbd是更传统的基准测试套件，其测试针对特定框架和模型进行了配置，涵盖了许多不同的机器学习应用程序。 同时，最近的深度学习框架侧重于比较给定模型和跨框架的数据集的性能。
而deepbench本身并不使用框架，而是使用低级库来评估不同设备的机器学习性能。就其本身而言，虽然它并不直接将框架/模型/应用程序性能与其他测试联系在一起，但它提供了代表供应商优化的数学操作和硬件性能的指标，每个产品的二进制文件都使用硬件供应商提供的库进行编译。
dawnbench则更加与众不同，与其说它是一个基准测试套件，不如说是对三个数据集（imagenet、cifar10和squad）的训练和推断结果进行类似于竞赛的报告，重点考量端对端的计算精确度和成本。
至于hpe dlbs，作为hpe深度学习指南的一部分，它主要以gpu为中心，坚持使用tensorflow、mxnet、pytorch和caffe类型框架，还包括tensorrt测试。虽然其具有良好的多测试批处理、日志记录、监控和报告功能，但它只输出纯粹的性能和时间指标，不涉及端对端的时间精度或成本。
从这些基准测试软件中可以看出，深度学习框架之间的差异很容易使测试结果变得毫无意义，从而影响我们对这些框架的研究。convnet-benchmark和pytorch的创始人soumith chintala指出，如果没有机器学习的背景，很难独立地验证深度学习基准测试的准确性和范围，不过mlperf测试项目似乎试图解决这个问题。
mlperf是由dawnbench等测试软件的设计者和工程师联合打造的全新高端基准测试套件，希望囊括fathom的跨域测试方法以及dawnbench对超过阈值精度模型的端对端计算时间考察。不过它目前正在处于alpha阶段，开发团队表示其尚不适合进行精确的硬件对比。
综合考虑之下，本次测试将不包含mlperf项目，而是使用deepbench、caffe2 docke、stanford dawn和hpe dlbs来进行。
deepbench训练测试之gemm和rnn
首先进行的是gemm测试，利用某些深度学习应用程序（deepspeech、speaker id和language modeling）中的内核进行gemm操作，测出的性能比在cublas中运行纯矩阵-矩阵乘法更有代表性。
测试的结果在意料之内，启用tensor core可以大幅提升性能。深入研究细节可以发现，tensor core对于特定类型的矩阵-矩阵乘法会有特别的影响。
通过深度学习应用程序拆分gemm测试，我们可以了解tensor core在理想和非理想情况下的表现。
speaker id gemm工作负载实际上只包含两个内核，其中10微秒的时间差意味着大约1 tflops的算力差异。
通过对语言模型内核的研究，可以了解tensor core在非理想情况下的性能。这些核矩阵的大小是m=512或1024，n=8或16，k=500000，虽然每个数在技术上都可以被8整除——这是满足张量核加速度的基本要求之一——但这些矩阵的形状与tensor core支持的16*16*16、32*8*16和8*32*16等基本wmma形状不太匹配。假如tensor core真正在独立的8x8x8级别上运行，那么运算8*8*8矩阵的性能也不会很好。
因此，tensor core无法高效的将这些非常不平衡的矩阵分解为n=8或16。而且，tensor core在deepspeech内核上的性能也出现异常：
从所有子项的平均成绩来看，这个浮点运算性能令人印象深刻。当矩阵适合于tensor core时，性能可以超过90tflops；相反如果二者无法契合，并正确的换位没有发挥作用，性能会低至<1tflops的水平。
对于deepbench rnn内核的测试，rnn类型之间没有明显的差异，但是在每种rnn类型中，如果将不同内核挨个进行对比判断，也可以看到与gemm中相同的趋势。
比较有趣的是，titan xp与titan v在未使用tensor core加速时的表现有很接近，titan xp的高频率为其性能起到了一定的帮助。
deepbench训练测试之convolutions
在卷积训练工作负载测试中，tensor core再次显着提高了性能。鉴于卷积层是图像识别和分类的基础，因而卷积运算是tensor core加速的最大潜在受益者之一。
从所有测试项的平均成绩可以看出，volta在启用了tensor core的fp16混合精度运算能力后性能再次取得了领先。不过与gemm不同，在fp32卷积上启用tensor core会导致明显的性能损失。
当计算涉及不匹配的张量尺寸时，标准精度模式遵循cudnn指定的最快前向算法（如winograd），而混合精度模式必须对所有内核使用隐式预计算gemm，这会造成两种混合精度模式的性能会出现下滑。
要符合tensor core加速的要求，输入和输出通道尺寸必须是8的倍数，输入、过滤和输出数据的类型必须是半精度。使用tensor core实现卷积加速要求张量采用nhwc格式，但大多数框架都希望采用nchw格式的张量。在这种情况下，输入通道不是8的倍数，但测试程序会自动填充以解决此问题。
需要注意的是，所有这些nchw内核都需要转换为nhwc。想要从tensor core中受益，需要正确的调整卷积格式，本次测试使用的是nvidia提供的标准库和makefile。nvidia指出，一旦进行加速卷积，它会消耗掉相当多的运行时间，这将会对fp32和fp16混合精度模式造成影响。
deepbench推理测试之gemm
数据精度方面，百度将deepbench gemm和卷积定义支持32bit累加的int8格式，以支持volta和pascal上的int8计算。
titan v和titan xp均拥有4倍于int32的int8性能，deepbench的int8推理测试正中pascal引入的dp4a矢量点积能力之下怀。volta同样拥有这一能力，在指令集中二者均显示为idp和idp4a。
对igemm来说，正如cutlass所示，dp4a是一项定制操作。因此除语言建模之外，int8的性能都非常之高。当然，与硬件不匹配的张量尺寸不适合tensor core加速，这一点与之前完全一样。
在完全连接（仿射）层中，每个节点都与前一层中的各节点相连接。对于一个典型的cnn来说，完全连接的层意味着可以结合所有提取的特征做出最终预测并对图像进行分类。这些测试结果数据也意味着大型且规则的矩阵可以在tensor core加速中获得更大的收益。
deepbench推理测试之convolutions
再次来到卷积测试环节，8位乘法/32位累加再次出现在int8推理中。
测试中最引人注目的是titan xp，在resnet、speaker id和vision项目中，titan xp表现出了强劲的int8吞吐量。
从内核方面来看，并没有发现这一现象的根源所在，猜测可能是由于pascal的dp4a库好驱动程序比volta更为成熟所致，亦或许是volta通过单独的int单元处理这些运算。
deepbench推理测试之rnn和sparse gemm
deepbench的最后一项推理测试是rnn和sparse gemm，虽然测试中可以选择fp16，但实际上它们都只支持fp32运算。
虽然rnn可能会有加速，但deepbench和nvidia目前仅支持单精度rnn推理。
nvidia caffe2测试之resnet50和imagenet
虽然内核和深度学习数学运算可能很有用，但实际应用中是使用真实数据集进行训练的。使用标准的ilsvrc 2012图片集，在imagenet上通过resnet50模型来训练和推断，可以展示更具参考的性能数据。
虽然fp16和tensor core有单独的开关，但titan v在启用和禁用tensor core的情况下运行fp16的性能是完全一样的。
只看原始吞吐量性能的话，titan v在所有批尺寸下都处于领先地位。凭借tensor core，titan v可处理的批尺寸达到甚至超过了64，而其他显卡即便有12 gb显存也无法与之相比。
不过只看原始吞吐量性能的问题在于，深度学习的实际性能从来没有这么简单。首先，许多模型可能会牺牲精度和训练时间以换取针对吞吐量的优化，如果模型需要较长的时间来收敛，那么每秒训练的峰值性能就没有参考意义了。
这些问题与使用fp16存储和tensor core的volta尤为相关，如果在实际应用中使用了有损缩放或单精度批量归一化，这在吞吐量性能中都是无法体现的。
hpe dlbs caffe2测试之resnet50和imagenet
接下来，我们看一下深度学习指南中的hpe dlbs。与通常的深度学习测试不同，hpe dlbs基本上只输出吞吐量和时间指标。
hpe dlbs的一大特色是支持nvidia caffe2测试使用的数据集，我们同样可以在imagenet上使用resnet50模型来训练和推断。但是由于二者的模型和实现不同，测出的吞吐量性能无法与nvidia caffe2直接进行比较。
在测试中，titan v无法支持某些特定的批尺寸，但总体趋势和之前的测试基本相同，fp16和tensor core提供了更高的吞吐量。不过遗憾的是，hpe dlbs caffe2测试似乎不支持int8推理。
hpe dlbs tensorrt测试之resnet50和imagenet
hpe dlbs的另一大特色是支持tensorrt（nvidia推理优化引擎）的基准测试功能， nvidia近年来已将tensorrt与新的深度学习功能（如int8/dp4a和tensor core的16位累加器模式）相结合以进行推理。
使用caffe模型，tensorrt可以根据需要调整模型，以便在给定的精度下进行推理。我们在titan x（maxwell）和titan xp（pascal）上运行了64、512和1024的批尺寸，在titan v运行了128、256和640的批尺寸。
titan xp的高int8性能在一定程度上印证了gemm/卷积性能，这两个工作负载似乎都在使用dp4a。不过雷锋网(公众号：雷锋网)并未了解到dp4a如何在titan v上实现，只知道它由volta指令集提供支持，且volta确实拥有独立的int32单元。
dawnbench测试之cifar10图像分类
就实际应用的性能而言，深度学习训练更适合用时间/准确性和成本来描述，而这两点分别对应dawnbench的两项子测试。对于使用cifar10的图像分类来说，这两项子测试为：
时间/准确性：训练cifar10数据集的图像分类模型，报告训练所需的时间，且要求测试集准确性至少为94%。
成本：在公共云基础架构上，计算达到94％或更高的测试集准确性所需的总时间，将所花费的时间（以小时为单位）乘以每小时实例的成本，以获得训练模型的总成本。
测试选用pytorch的cifar10训练实现中最快的两个，其中一个基于resnet34，是为了在nvidia geforce gtx 1080 ti上运行，而第二个基于resnet18，是为了在单个tesla v100上运行。这些都是dawnbench中最近的热门测试，可以认为它们是相当符合现代的项目，同时cifar10也不是一个非常密集的数据集。
cifar10的小型图像数据集运行良好，第一个训练实现是在单个gtx 1080 ti上运行，需要35分37秒才能训练到94％的准确性，而在第二个训练实现中，titan v只用了5分41秒就完成了94%的目标。
顺带一提，虽然titan v在第一个训练实现中不会使用tensor core，但凭借相对于pascal的一般改进，titan v在这个测试中的速度依然比titan xp快20％左右，同时系统峰值功耗也下降了大约80w。
结语
tensor core是titan v的重要组成部分，本文的目的也是尽可能的了解volta，所以测试着重考察了tensor core加速。
本次测试还有许多没有涉及到的各种其他测试和套件，它们普遍在设计上有太多欠缺。事实证明，无论从整体还是局部来看，都不可能找到一个负载贴合当下实际、提供端到端指标、覆盖多个机器学习域、支持tensor core和混合精度，最重要的是易于非专业开发者使用的深度学习套件。
即便是参考价值较大的dawnbench，设计本意也并非提供一个通用的基准，而是为方便研究者和开发者创建他们自己的实现。dawnbench的深度学习框架仍然需要修改以作为一款有效的基准测试使用，但重新配置一个与volta兼容的混合精度模型就不是一件可以轻松完成的事情。
这实际上与titan v自身相关，tensor core和混合精度需要专门的开发来适配支持，只能在特定情况下明显提升性能，而且即便用户的代码和数据集可以与tensor core良好的匹配，到头来神经网络处理也会受到来自传统alu的限制。
而对于主流消费者群体来说，tensor core的发展对他们意味着什么？
最新的turing架构证明，tensor core在游戏卡领域同样可以有所应用。rtx光线追踪技术就是使用tensor core对图像进行降噪以弥补光线数量的限制。nvidia唯一需要考虑的问题是，这些玩家是否乐意为这些tensor core买单。
对于任何考虑购买titan v以满足计算需求的人来说，投资titan v就意味着投资混合精度深度学习模型和基于wmma的hpc gemm加速。在cudnn、cublas以及早期dp4a和fp16*2混合精度计算的背景下，tensor core是试图用可编程硬件来实现深度学习加速的一种自然进化。
可以肯定的是，titan v绝对代表了nvidia对未来gpgpu的愿望。
2021中国汽车供应链峰会暨第六届铃轩奖盛典圆满落幕
数据网格是什么意思？为什么要考虑实现数据网格？
Anritsu MS8604A数字移动无线发射机测试仪孟姐/
结合从业者Michail Katkoff的数据分享以及《一起来捉妖》的游戏特色，来分析该品类的现状和发展趋势
GCAN以太网can网关如何修改本地IP？
NVIDIA Tensor Core深度学习核心解析及跑分测试
debian软件包管理详解
ADI公司推出集成式隔离电源控制器系列
国民经济的基础性行业之一---减速机行业
亚马逊涉嫌未能保护工人免受COVID-19的侵害而被起诉
无位置反馈比例阀放大器RT-PVDA
AMD官方澄清RX5700XT50周年纪念版显卡限地区发售的情况
全球领先系统制造商推出 NVIDIA AI 就绪型服务器，为企业大幅提升生成式 AI 性能
ds18b20如何读取温度？温度数据怎么换算？
Linux操作系统的系统性能监控工具syastat的详解
首创后置三摄+三摄万物：华为联合京东这波海报666
物联网安全之Unicorn框架初探
高频闭锁方向保护的动作条件和闭锁条件是什么？
七月新机红米Pro2，屏幕占比史无前例，价格适中1999
华为Mate50“没电也能打电话”