这个问题表明它是在机器学习的背景下提出的,因此重点是单精度计算,特别是使用 IEEE-754binary32
格式。Asker 表示,NVIDIA GPU 是一个令人感兴趣的平台。我将专注于使用 CUDA 使用这些 GPU,因为我不熟悉 CUDA 的 Python 绑定。
当谈到 FLOPS 时,除了简单的加法和乘法之外,还有各种关于如何计算它们的思想流派。GPU 例如计算软件中的除法和平方根。识别浮点指令并对它们进行计数就不会那么模棱两可了,这就是我在这里要做的。请注意,并非所有浮点指令都将以相同的吞吐量执行,这也可能取决于 GPU 架构。有关指令吞吐量的一些相关信息可以在 CUDA 编程指南中找到。
从图灵架构(计算能力 7.5)开始,GPU 包含一条指令MUFU.TANH
,用于计算精度约为 16 位的单精度双曲正切。多功能单元 (MUFU) 支持的单精度函数通常通过存储在 ROM 中的表中的二次插值计算。我能说的最好的,MUFU.TANH
是在虚拟汇编语言 PTX 的级别上公开,但不是(从 CUDA 11.2 开始)作为设备函数内在函数。
但是鉴于功能是在 PTX 级别公开的,我们可以通过一行内联汇编轻松创建自己的内在函数:
// Compute hyperbolic tangent for >= sm75. maxulperr = 133.95290, maxrelerr = 1.1126e-5
__forceinline__ __device__ float __tanhf (float a)
{
asm ("tanh.approx.f32 %0,%1; \n\t" : "=f"(a) : "f"(a));
return a;
}
MUFU.EX2
在计算能力 < 7.5 的旧 GPU 架构上,我们可以通过代数变换和使用机器指令和来实现具有非常匹配特征的内在函数MUFU.RCP
,它们分别计算指数基数 2 和倒数。对于幅度较小的参数,我们可以使用 tanh(x) = x 并通过实验确定两个近似值之间的良好切换点。
// like copysignf(); when first argument is known to be positive
__forceinline__ __device__ float copysignf_pos (float a, float b)
{
return __int_as_float (__float_as_int (a) | (__float_as_int (b) & 0x80000000));
}
// Compute hyperbolic tangent for < sm_75. maxulperr = 108.82848, maxrelerr = 9.3450e-6
__forceinline__ __device__ float __tanhf (float a)
{
const float L2E = 1.442695041f;
float e, r, s, t, d;
s = fabsf (a);
t = -L2E * 2.0f * s;
asm ("ex2.approx.ftz.f32 %0,%1;\n\t" : "=f"(e) : "f"(t));
d = e + 1.0f;
asm ("rcp.approx.ftz.f32 %0,%1;\n\t" : "=f"(r) : "f"(d));
r = fmaf (e, -r, r);
if (s < 4.997253418e-3f) r = a;
if (!isnan (a)) r = copysignf_pos (r, a);
return r;
}
使用 CUDA 11.2 为sm_70
目标编译此代码,然后使用cuobjdump --dump-sass
显示八个浮点指令反汇编二进制文件。我们还可以看到生成的机器代码 (SASS) 是无分支的。
如果我们想要一个具有完全单精度精度的双曲正切,我们可以对量级较小的参数使用极小极大多项式近似,同时对量级较大的参数使用代数变换MUFU.EX2
和机器指令。MUFU.RCP
超过一定幅度的参数,结果将为±1。
// Compute hyperbolic tangent. maxulperr = 1.81484, maxrelerr = 1.9547e-7
__forceinline__ __device__ float my_tanhf (float a)
{
const float L2E = 1.442695041f;
float p, s, t, r;
t = fabsf (a);
if (t >= 307.0f/512.0f) { // 0.599609375
r = L2E * 2.0f * t;
asm ("ex2.approx.ftz.f32 %0,%1;\n\t" : "=f"(r) : "f"(r));
r = 1.0f + r;
asm ("rcp.approx.ftz.f32 %0,%1;\n\t" : "=f"(r) : "f"(r));
r = fmaf (r, -2.0f, 1.0f);
if (t >= 9.03125f) r = 1.0f;
r = copysignf_pos (r, a);
} else {
s = a * a;
p = 1.57394409e-2f; // 0x1.01e000p-6
p = fmaf (p, s, -5.23025580e-2f); // -0x1.ac766ap-5
p = fmaf (p, s, 1.33152470e-1f); // 0x1.10b23ep-3
p = fmaf (p, s, -3.33327681e-1f); // -0x1.5553dap-2
p = fmaf (p, s, 0.0f);
r = fmaf (p, a, a);
}
return r;
}
此代码包含一个数据相关分支,查看 CUDA 11.2 为sm75
目标生成的机器代码表明该分支被保留。这意味着一般来说,在所有活动线程中,一些将遵循分支的一侧,而其余的将遵循分支的另一侧,需要后续同步。因此,为了获得有关所需计算工作量的现实想法,我们需要结合两个执行路径的浮点指令计数。这产生了 13 条浮点指令。
上面代码注释中的错误界限是通过对所有可能的单精度参数的详尽测试建立的。