提问人：Agade 提问时间：3/30/2015 最后编辑：phuclvAgade 更新时间：9/10/2020 访问量：5030
在 GPU 上使用 2 个 FP32 模拟 FP64
Emulating FP64 with 2 FP32 on a GPU
问：
如果用两个单精度浮点来模拟双精度浮点，性能会是什么样子，能做得好吗？
目前，英伟达对支持双精度的特斯拉卡收取相当高的费用，这使您可以获得单精度性能的三分之一（值得注意的例外：Titan/Titan Black）。
如果使用具有双精度的 Geforce GPU，并使用 2 个单精度浮点数模拟双精度，性能会如何？
答：
14赞
njuffa
3/31/2015
#1
您可以通过计算实现每个双浮点操作所需的操作数来粗略估计性能。您可能希望检查二进制代码以获得准确的计数。我在下面展示了一个双浮点乘法，它充分利用了 GPU 上的 FMA（融合乘加）支持。对于双浮点数加法代码，我会向你指出 Andrew Thall 的一篇论文，因为我现在没有时间对此进行编码。从之前的分析来看，我相信论文中给出的加法代码是正确的，并且它避免了更快但不太准确的实现中的常见陷阱（当操作数的大小在两倍以内时，这些错误会失去准确性）。floatcuobjdump --dump-sass
如果您是注册的 CUDA 开发人员，您可以从 NVIDIA 的开发人员网站（在 https://developer.nvidia.com 登录）下载 double-double 代码，该网站在 BSD 许可下，并相对较快地将其重新编写为 double-float 代码。NVIDIA 的双双代码支持运算加法、减法、除法、平方根和倒数平方根。
如您所见，下面的乘法需要 8 条指令;一元否定被吸收到 FMA 中。添加大约需要 20 条指令。然而，双浮子操作的指令序列也需要临时变量，这会增加寄存器压力并降低占用率。因此，一个相当保守的估计可能是，双浮点数算术的吞吐量是本机算术的 1/20。您可以在与您相关的上下文（即您的用例）中轻松衡量这一点。floatfloatfloat
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;
}
请注意，在各种应用中，可能不需要完整的双浮点数算术。取而代之的是，人们可以使用计算，并通过误差补偿技术进行增强，其中最古老的技术之一是卡汉求和。我最近在 NVIDIA 开发者论坛上发表的一篇文章中简要概述了有关此类方法的易用文献。在上面的评论中，Robert Crovella还提到了Scott LeGrand在GTC 2015上的演讲，我还没有时间去看。float
至于精度，double-float 的表示精度为 49 （24+24+1） 位，而 IEEE-755 的表示精度为 53 位。然而，对于小幅度的操作数，双浮点数无法保持这种精度，因为尾部可能会变成非正态或零。当非正规支持打开时，保证 2-101 <= |x|< 2128.默认情况下，CUDA 工具链中对架构 >= sm_20 的非正常支持处于打开状态，这意味着当前发布的版本 CUDA 7.0 支持的所有架构。doublefloat
与对 IEEE-754 数据的运算相反，双浮点运算未正确舍入。对于上面的双浮点数乘法，使用 20 亿个随机测试用例（所有源操作数和结果都在上述范围内），我观察到相对误差的上限为 1.42e-14。我没有双浮点数加法的数据，但它的误差范围应该是相似的。double
评论