一、前言

FP8 是 FP16 的衍生产物，它包含两种编码格式 E4M3 与 E5M2。对于 E4M3 而言，其包含 4 比特指数、3 比特底数、以及 1 比特符号位。

E5M2 同理包含 5 比特指数位、3 比特底数、1 比特符号。在本文中，我们称指数部分为 exponent，底数部分为 mantissa。

下图展示了 FP32, FP16, FP8 的格式对比：

使用 FP8 数据格式，能够有效地以在推理时提升效率，同时也可以使用 FP8 数据格式直接进行训练，从而到达 16bit 训练的精度。在后续的内容中 Nvidia 的开发者在常见模型中使用 FP8 数据格式进行实验，其实验结果可以媲美 FP16 的精度。

二、FP8 数据格式

在本节中，我们将向你介绍FP8的表示方法，以及使用FP8表示神经网络数据时的细节问题。

2.1 浮点数表示法：

在计算机中我们使用符号位、指数、底数三部分表示一个浮点数。符号位只占用1bit，用来表达数的正负，0-表示整数，1-表示负数。这也意味着在浮点数中0有两种表达方式：

二进制的 (0 0000 000) 表示 FP8 中的正数 0

二进制的 (1 0000 000) 表示 FP8 中的负数 0

接下来我们讨论浮点表示中的指数部分：

除符号位外浮点数具有一个指数部分，这也是其区别于定点数的地方。在 FP8 E4M3 格式中，我们具有4个比特用于表示指数。在(0 1000 000) 中，指数部分的 1000 为十进制的 8，其对应的十进制数为：2^(8 - 7) * (1) = 2

在浮点表示法当中，指数部分总是会减去一个偏移量，对于FP8 E4M3 而言，这个偏移量为-7，这使得指数的表示范围为[-7, 8]。对于 FP8 E5M2 而言，指数偏移量为 -15，指数表示范围为[-15, 16]。

特别地，我们规定当指数部分全为1时的浮点值为无穷大，这一规定也适用于 FP16 与 FP32，是IEEE754标准的一部分。

二进制的 (0 11111 10) 在 FP8 E5M2 格式中表示正无穷大

二进制的 (1 11111 00) 在 FP8 E5M2 格式中表示负无穷大

FP8 E4M3 不遵循 IEEE754 标准，其在指数全为1时仍然可以表示有效数字，当且仅当指数与底数部分全为1时，其表示无穷大。

二进制的 (0 1111 110) 在 FP8 E4M3 格式中表示 2^(15-7) * (1+1/2 + 1/4) = 448

二进制的 (1 1111 100) 在 FP8 E4M3 格式中表示 -2^(15-7) * (1+1/2) = 384

二进制的 (1 1111 111) 在 FP8 E4M3 格式中表示负无穷大

最后我们讨论浮点数中的底数部分：

在浮点数中，底数从高位到低位，分别表示2的负k次幂；对于E4M3格式，我们使用3个比特表示底数，其分别对应2的负1, 2, 3次幂（见上图2）。对于E5M2格式，我们使用2个比特表示底数，分别对应2的负1, 2幂。浮点数的底数值转换成十进制时，我们需要将所有底数位的数字相加，再额外加1。上图2中，左边的底数部分为: 1 + 1/2 + 1/8 = 1.625，右侧的底数部分为 1+0.25 = 1.25。

非规格化浮点数：

当一个浮点数的指数部分全为0时，其是一个非规格化浮点数，此时其底数部分不再额外加1。

(0 0000 000) 表示 2^(0-7) * 0，而非 2^(0-7) * 1

(0 0000 001) 表示 2^(0-7) * 1/8

2.2 数的浮点量化

由于指数部分的存在，浮点量化与INT量化表现出了不一样的性质，浮点量化的步长是"可变"的。从图像上来看，浮点量化器的量化步长随着指数部分的变大而变大。在本文中，我们以红色表示FP8 E4M3的量化情况，蓝色表示FP8 E5M2的量化情况，而绿色表示INT8的量化情况：

FP8 E4M3的表示范围为[-448, 448]，INT8的表示范围为[-128, 127]；在图3中的实验中，我们进行INT8量化时额外地添加了缩放因子scale = 4.0，从而保证在[0, 500]范围上三种方案都能够比较良好地完成表示。

图上的例子说明，INT的量化步长是均匀的，总是以一定的步长完成量化，这是一种均匀的量化。而浮点的量化则是非均匀的，随着数值增大，其步长也在逐渐变大。且E5M2的步长变化较E4M3而言更加明显。从另一个角度出发，量化的误差总是与步长正相关的，因此FP8浮点量化相比于INT8而言，对于小数来说会更加精确，但对于大数则更不精确。

2.3 FP32 到 FP8 的数据格式转换

FP8 E4M3 的表示范围只有[-448, 448]，对于一些应用来说，这是无法完成表示的。因此我们也额外地在转换过程中引入缩放因子 scale。FP32 向 FP8 的转换过程可以表示为先除以尺度因子得到中间结果 Unscaled FP32，再由中间结果完成 FP32 到 FP8 的格式转换。

Unscaled FP32 = FP32 / scale
FP8 = Convert(Unscaled FP32)

下面我们讨论Convert函数的执行过程，我们将讨论三种情况：

当 Unscaled FP32 数据已经超出 FP8 的表示范围，即 Unscaled FP32 的幅值大于 448，那么直接进行截断，此时为浮点上溢出。
当 Unscaled FP32 数据范围在 FP8 的表示范围内，且幅值大于 FP8 能够表达的最小值，此时需要移去多余的底数位，并对底数进行四舍五入。
当 Unscaled FP32 数据小于 FP8 能够表达的最小值，此时浮点下溢出，只需判断能否四舍五入为 (0 0000 001)，若不能则直接为0。

以程序实现，代码如下：

union FPConvertHelper {
    float value;
    uint32_t data;
};

template<typename Dtype, typename Stype, typename Otype>
__device__ __inline__
float QuantizeScalarFloating(
    const Dtype value, const Stype scale, const Otype offset,
    const int exponent, const int mantissa,
    const float clip_min, const float clip_max, 
    const Rounding rounding){
    /**
     * PPQ Quantization Function implementation.
     * This function convert an float value to low-precision float
     */
    FPConvertHelper helper; FPConvertHelper rounding_helper;
    helper.value = static_cast<float>(value) / scale;
    // Following code will Split float32 into sign, exp, mantissa
    /* IEEE 754 Standard: 1 bit sign, 8 bit exponent, 23 bit mantissa */

    /* In binary 10000000 00000000 00000000 00000000 = 0x80000000 in Hex */
    /* In binary 01111111 10000000 00000000 00000000 = 0x7F800000 in Hex */
    /* In binary 00000000 01111111 11111111 11111111 = 0x007FFFFF in Hex */

    /* Tool: https://www.h-schmidt.net/FloatConverter/IEEE754.html */

    uint32_t fp32_sign    = helper.data & 0x80000000;
    int32_t fp32_exp      = helper.data & 0x7F800000;
    int32_t fp32_mantissa = helper.data & 0x007FFFFF;
    int32_t exponent_min  = -(1 << (exponent - 1)) + mantissa;
    int32_t exponent_max  = (1 << (exponent - 1));

    // Float Overflow.
    if (value > clip_max) return clip_max;
    if (value < clip_min) return clip_min;

    // Following code will process Float underflow
    /* Float underflow means fp32_exp is smaller than exponent_min          */
    /* Where exponent_min is the minimum exponent value of quantized float. */
    /* For FP8 E4M3, the minimum exponent value should be -9.               */
 if (((fp32_exp >> 23) - 127) < exponent_min){
        if (((fp32_exp >> 23) - 127) == (exponent_min - 1)){
            // there is a chance to round
            rounding_helper.data = (fp32_mantissa & 0x007FFFFF) + 0x3F800000;
            if (_round2int(rounding_helper.value - 1, rounding)) {
                helper.data = fp32_sign + ((exponent_min + 127) << 23) + (1 << (23 - mantissa));
                return helper.value;
            }
        }
        return 0.0f;
 }
    if ((fp32_exp >> 23) - 127 > exponent_max){
        if (fp32_sign) return clip_min;
        else return clip_max;
    }

    /* high precision mantissa convert to low precision mantissa requires rounding                         */
    /* Here we apply a tricky method to round mantissa:                                                    */
    /* We create another float, which sign = 0, exponent = 127, mantissa = fp32_mantissa << (23 - mantissa) */
    /* Then we directly round this float to int, result here is what we want, you can prove it by yourself */
    rounding_helper.data = ((fp32_mantissa << (mantissa)) & 0x007FFFFF) + 0x3F800000;
    uint32_t round_bit = _round2int(rounding_helper.value - 1, rounding);

    // process mantissa
    fp32_mantissa = ((fp32_mantissa >> (23 - mantissa)) + round_bit) << (23 - mantissa);
    helper.data = fp32_sign + fp32_mantissa + fp32_exp;

    return CLIP<float>(helper.value + offset, clip_min, clip_max);
}

三、浮点量化误差分析

FP8 具有更大的表示范围，但在一定范围内，其表示精度相较 INT8 更差。为了阐明这一问题，我们从正太分布中随机抽样 1000 万个数字，分别使用 FP8 E4M3, FP8 E5M2, INT8 完成量化。

在三者的量化中，我们都应用缩放因子来调整量化效果。

通过不断调整缩放因子的大小，我们可以得到量化的误差随缩放因子的变化情况，这反映了一些有意义的问题：

我们认为在选取了合适的缩放因子时，INT8的量化精度高于FP8，两者之间的误差几乎相差一个数量级。这是INT8量化的优势，它更加精确。FP8将提供更好的宽容性，在scale的选取上不需要花费太多力气，FP8的量化对scale的选择不敏感。也就是说FP8的场景下，你几乎可以不做 calibration 就完成网络的量化。

在之前的视频中我们已经阐述过这样的观点：量化误差是一个关于量化步长的二次函数，随着scale的增大其误差将很快发散。FP8的量化同样遵循这个原理，但其表示范围太过宽泛，它的误差曲线是一个开口非常大的二次函数，从图上看它几乎是平的。只有当scale选取非常不合理时，FP8的误差才会呈现二次曲线的性质。

在图6中，我们进一步扩大scale的选取范围，可以看到对于正太分布而言，scale=2 时 FP8 E4M3 仍然有较高的量化精度，而此时int8的量化早已完全失效。同时FP8量化误差曲线呈现明显的周期性，读者可以自行探索这一现象发生的原因。

在这里，我们不想进一步分析FP8的量化误差公式，这些公式非常难算且没有太多意义。在后续更新的视频内容中我们可能会对这一部分内容有更深入的讨论。我们只给出一个有参考意义的结果。下表反映了对于正太总体而言，以最优的scale完成量化，FP8与INT8的量化误差期望。其中量化误差的衡量方式为量化噪声的能量除以信号自身的能量，可以看到INT8的误差期望是要远小于FP8的。

四、实验部分

在这里我们给出两部分实验结果，第一部分来自于 Nvidia 的论文，在论文中研究者使用了训练中量化的技术，从而训练出了一批适合于FP8的网络模型，并与直接使用FP16完成训练的模型进行了精度对比，实验结果如下：

浮点量化的训练过程是简单的，由于浮点量化对scale并不敏感，我们可以在训练时直接省去对scale的讨论，将所有缩放因子设置为1.0，只需保证网络中的激活与权重在训练过程中不会超出FP8的表示范围，也即[-448, 448]。相比FP32，乐观地讲我们只需要在网络中加入一些额外的clip算子就可以完成FP8的QAT训练。一个额外的细节是：在训练时，如果我们发现网络中的数值或梯度出现了inf，我们可以降低网络的学习速率来保证稳定。

第二部分的实验结果来自我们的量化工具

在这部分实验中我们不加训练地直接量化分类网络，设置所有缩放因子为1.0，同时以前文提及的数据转换函数执行FP8的量化模拟，实验结果如下：

在两部分的实验中，我们均只量化卷积和矩阵乘的输入和权重，网络中的其他值保持为FP32精度，Bias也保持为FP32精度。

五、结论

FP8的量化并不精确
FP8的量化具有良好的宽容度，我们期待他在QAT中取得更好的表现
FP8良好的宽容度可以量化一些奇怪的网络，例如Effcientnet。

总的来看，FP8 更加适合训练，这是一项很大的优点，其良好的宽容度让我们可以期待 FP8 QAT 有着更好的稳定性，用户也可以借 FP8 省去非常麻烦的 Calibration 过程。FP8 的量化是不精确的，另一方面 FP8 的运算器设计依然会比 INT8 的更加复杂，它依然不适合端侧芯片的部署，更加适合于 GPU。

▎参考：

· https://github.com/openppl-public/ppq/tree/master/ppq
· https://github.com/openppl-public/ppq/pull/274
· https://www.graphcore.ai/posts/graphcore-and-amd-propose-8-bit-fp-ai-standard-with-qualcomm-support
· 本文部分内容翻译自：https://arxiv.org/pdf/2209.05433.pdf

作者：张志
文章来源：OpenPPL

推荐阅读

更多嵌入式AI干货请关注嵌入式AI专栏。欢迎添加极术小姐姐微信（id:aijishu20)加入技术交流群，请备注研究方向。

FP8 量化：原理、实现与误差分析