在 GPU 上用 2 个 FP32 模拟 FP64

Question

如果用两个单精度浮点来模拟双精度浮点，性能会怎样，能不能做好？

目前，Nvidia 对启用双精度的 Tesla 卡收取相当高的费用，这使您可以获得单精度性能的三分之一（值得注意的例外 Titan/Titan 黑色）。

如果使用具有 gimped 双精度的 Geforce GPU 并使用 2 个单精度浮点数模拟双精度，性能会怎样？

Answer 1

您可以通过计算实现每个双浮点运算所需的 float 运算次数来粗略估计性能。您可能希望使用 cuobjdump --dump-sass 检查二进制代码以获得准确的计数。我在下面展示了一个双浮点乘法，它充分利用了 GPU 上的 FMA（融合乘加）支持。对于双浮点数加法代码，我会向您指出 a paper by Andrew Thall，因为我现在没有时间对此进行编码。根据之前的分析，我相信论文中给出的加法代码是正确的，并且它避免了更快但不太准确的实现中的常见陷阱（当操作数的大小在两倍以内时会失去准确性）。

如果您是注册的 CUDA 开发人员，您可以从 NVIDIA 的开发人员网站（登录 https://developer.nvidia.com）下载双精度代码，该代码在 BSD 许可下，并相对快速地将其重写为双浮点代码。 NVIDIA双双码支持加减除除平方根倒数平方根运算

如您所见，下面的乘法运算需要 8 float 条指令；一元否定被吸收到 FMA 中。添加需要大约 20 float 条指令。但是，双浮点运算的指令序列也需要临时变量，这会增加寄存器压力并降低占用率。因此，一个合理的保守估计可能是双浮点运算的吞吐量是原生 float 运算的 1/20。您可以在与您相关的上下文（即您的用例）中轻松地自行衡量这一点。

typedef float2 dblfloat;  // .y = head, .x = tail

__host__ __device__ __forceinline__ 
dblfloat mul_dblfloat (dblfloat x, dblfloat y)
{
    dblfloat t, z;
    float sum;
    t.y = x.y * y.y;
    t.x = fmaf (x.y, y.y, -t.y);
    t.x = fmaf (x.x, y.x, t.x);
    t.x = fmaf (x.y, y.x, t.x);
    t.x = fmaf (x.x, y.y, t.x);
    /* normalize result */
    sum = t.y + t.x;
    z.x = (t.y - sum) + t.x;
    z.y = sum;
    return z;
}

请注意，在各种应用程序中，可能不需要完整的双浮点运算。相反，人们可以使用 float 计算，并通过误差补偿技术进行增强，其中最古老的技术之一是 Kahan summation. I gave a brief overview of easily available literature on such methods in a recent posting in the NVIDIA developer forums. In the comments above, Robert Crovella also pointed to a GTC 2015 talk by Scott LeGrand，我还没有时间检查它。

至于精度，double-float 的表示精度为 49 (24+24+1) 位，而 IEEE-755 double 提供了 53 位。但是，对于数量级较小的操作数，双浮点数无法保持这种精度，因为尾部可能变为非正规或零。当打开非正规支持时，49 位精度保证 2^-101 <= |x| < 2¹²⁸。对于架构 >= sm_20，默认情况下在 CUDA 工具链中打开对 float 的非规范化支持，这意味着当前发布的版本 CUDA 7.0 支持所有架构。

与对 IEEE-754 double 数据的操作相反，双浮点数操作未正确舍入。对于上面的双浮点乘法，使用 20 亿个随机测试用例（所有源操作数和结果都在上述范围内），我观察到相对误差的上限为 1.42e-14。我没有双浮点数加法的数据，但它的误差界限应该是相似的。

在 GPU 上用 2 个 FP32 模拟 FP64

Emulating FP64 with 2 FP32 on a GPU

floating-point

cuda

double

double-double-arithmetic