【问题笔记】在NVIDIA Tesla V100上用非tensor core测试，半精度并不如单精度，是真的么？

有同学问：NVIDIA手册上说Tesla P100用sp做半精度，理论计算能力是单精度的一倍，我在Tesla V100上用非tensor core测试，半精度并不如单精度。不知道是p100是不是也这样？

对于这个问题我们好好唠唠:

备注/经验分享：

Tesla V100的半精度有两部分。 一部分是通用的，另外一部分是专用的。 通用的是SP，用户可以自行写代码进行任何涉及半精度的数据类型的运算。专用的是Tensor Core，只能进行矩阵乘法（图灵也是这样）。而在使用通用的半精度的时候，必须使用打包类型（2个half在一起）进行2个2个的运算。

这种才能100%。

NV有过文章：

The half type is used to represent FP16 values on the device. As with Kepler and Maxwell, FP16 storage can be used to reduce the required memory footprint and bandwidth compared to FP32 or FP64 storage. Pascal also adds support for native FP16 instructions. Peak FP16 throughput is attained by using a paired operation to perform two FP16 instructions per core simultaneously. To be eligible for the paired operation the operands must be stored in a half2 vector type. GP100 and GP104 provide different FP16 throughputs. GP100, designed with training deep neural networks in mind, provides FP16 throughput up to 2x that of FP32 arithmetic.

这段是CUDA自带的文章的。必须使用half2类型进行运算才能双倍。这是其一（这段来自Pascal的Tuning Guide。GP100只有通用的200%的FP16）

这里是CUDA 8当初的介绍，如何使用half2类型，请注意，和常规的直接使用加号（+）之类的运算不同，half2类型需要很多的intrinsic functions。

请用户直接参考此文，和打开cuda_fp16.h进行观察。需要注意的是，CUDA 9.2的这里稍微有些改变，但是改变不大。打开该头文件自己看即可（在你的CUDA 9.2的安装目录），注意。因为GP100这种卡，以及GV100和后续的图灵的通用计算，当使用FP16的时候必须2个一起上。如果用户只是用了1个half类型，会等于舍弃掉另外一半的计算性能。只有50%的200%的性能。也就是此时性能将和普通float一样（甚至一定条件下会不如普通float），为何？因为还有pack/unpack成本。所以建议在通用计算的时候，总是使用half2类型。

需要注意的是，如果能利用库建议利用，cublas之类的很多计算可以上半精度的。

大家也可以参考我们在论坛上的一个帖子：

一个寄存器能不能被两个half占用，比如__half a[2]={0， 0}是不是占一个寄存器 （点击阅读原文，访问这个帖子）