7

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

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

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

4

1 回答 1

12

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

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

可以看到，下面的乘法需要8float条指令；一元否定被吸收到 FMA 中。添加需要大约 20float条指令。但是，双浮点操作的指令序列也需要临时变量，这会增加寄存器压力并降低占用率。因此，一个合理保守的估计可能是双浮点算术的吞吐量是本地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 求和。在 NVIDIA 开发者论坛最近的一篇文章中，我简要概述了有关此类方法的容易获得的文献。在上面的评论中，Robert Crovella 还指出了Scott LeGrand 的 GTC 2015 演讲，我还没有时间查看。

至于精度，与double提供 53 位的 IEEE-755 相比，双浮点具有 49 (24+24+1) 位的表示精度。然而，双浮点数不能为数量级小的操作数保持这种精度，因为尾部可能变为非正规或零。打开非规范支持时，2 -101 <= |x|保证 49 位精度 < 2 128。默认情况下，在架构 >= sm_20float的 CUDA 工具链中启用非规范支持，这意味着当前发布的版本 CUDA 7.0 支持的所有架构。

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

于 2015-03-30T16:54:49.823 回答