tanh 需要多少 FLOP？

Question

我想计算 LeNet-5 (paper) 的每一层需要多少次失败。有些论文总共给出了其他架构的 FLOP（1、2、3）但是，这些论文没有详细说明如何计算 FLOP 的数量，我也不知道需要多少 FLOP非线性激活函数。例如，计算tanh(x)需要多少次FLOP？

我想这将是实现，也可能是特定于硬件的。但是，我主要对获得一个数量级感兴趣。我们是在谈论 10 FLOPs 吗？ 100 次翻牌？ 1000 次失败？因此，选择您想要的任何架构/实现作为答案。 （尽管我很欣赏接近“常见”设置的答案，例如 Intel i5 / nvidia GPU / Tensorflow）

Answer 1

如果我们查看 tanh(x) 的 glibc 实现，我们会看到：

1. 对于大于 22.0 和双精度的 x 值，可以安全地假定 tanh(x) 为 1.0，因此几乎没有成本。
2. 对于非常小的x，（假设是x<2^(-55)）另一个廉价的近似值是可能的：tanh(x)=x(1+x)，因此只需要两个浮点运算。
3. 对于beetween中的值，可以重写tanh(x)=(1-exp(-2x))/(1+exp(-2x))。不过一定要准确，因为1-exp(t)对于小t值由于失去显着性是非常有问题的，所以用expm(x)=exp(x)-1计算tanh(x)=-expm1(-2x)/(expm1(-2x)+2)。

所以基本上，最坏的情况大约是 expm1 所需操作次数的 2 倍，这是一个相当复杂的函数。最好的方法可能只是将计算 tanh(x) 所需的时间与两个双精度数的简单相乘所需的时间进行比较。

我在英特尔处理器上的（草率）实验产生了以下结果，这给出了一个粗略的想法：

因此对于非常小的数字和 >22 的数字，几乎没有成本，对于高达 0.1 的数字，我们支付 6 FLOPS，然后每个 tanh-caclulation 的成本上升到大约 20 FLOPS。

关键要点：计算 tanh(x) 的成本取决于参数 x，最大成本在 10 到 100 次 FLOP 之间。

---

有一个名为 F2XM1 的 Intel 指令为 -1.0<x<1.0 计算 2^x-1，至少在某些范围内可用于计算 tanh。但是，如果要相信agner's tables，则此操作的成本约为 60 FLOP。

---

另一个问题是向量化——据我所知，正常的 glibc 实现没有向量化。因此，如果您的程序使用矢量化并且必须使用未矢量化的tanh 实现，它将进一步减慢程序的速度。为此，英特尔编译器具有 mkl 库，其中 vectorizes tanh 等等。

正如您在表格中看到的，每次操作的最大成本约为 10 个时钟（浮动操作的成本约为 1 个时钟）。

---

我想你有机会通过使用-ffast-math 编译器选项赢得一些 FLOP，这会导致程序更快但不太精确（这是 Cuda 或 c/c++ 的选项，不确定是否可以这样做对于 python/numpy)。

---

为图生成数据的 c++ 代码（使用 g++ -std=c++11 -O2 编译）。它的目的不是给出确切的数字，而是对成本的第一印象：

#include <chrono>
#include <iostream>
#include <vector>
#include <math.h>

int main(){
   const std::vector<double> starts={1e-30, 1e-18, 1e-16, 1e-10, 1e-5, 1e-2, 1e-1, 0.5, 0.7, 0.9, 1.0, 2.0, 10, 20, 23, 100,1e3, 1e4};
   const double FACTOR=1.0+1e-11;
   const size_t ITER=100000000; 


   //warm-up:
   double res=1.0;
      for(size_t i=0;i<4*ITER;i++){
      res*=FACTOR;
   }
   //overhead:
   auto begin = std::chrono::high_resolution_clock::now();
   for(size_t i=0;i<ITER;i++){
      res*=FACTOR;
   }
   auto end = std::chrono::high_resolution_clock::now();
   auto overhead=std::chrono::duration_cast<std::chrono::nanoseconds>(end-begin).count(); 
   //std::cout<<"overhead: "<<overhead<<"\n";


   //experiments:
   for(auto start : starts){
       begin=std::chrono::high_resolution_clock::now();
       for(size_t i=0;i<ITER;i++){
           res*=tanh(start);
           start*=FACTOR;
       }
       auto end = std::chrono::high_resolution_clock::now();
       auto time_needed=std::chrono::duration_cast<std::chrono::nanoseconds>(end-begin).count();
       std::cout<<start<<" "<<time_needed/overhead<<"\n"; 
   }

   //overhead check:
   begin = std::chrono::high_resolution_clock::now();
   for(size_t i=0;i<ITER;i++){
      res*=FACTOR;
   }
   end = std::chrono::high_resolution_clock::now();
   auto overhead_new=std::chrono::duration_cast<std::chrono::nanoseconds>(end-begin).count(); 
   std::cerr<<"overhead check: "<<overhead/overhead_new<<"\n";
   std::cerr<<res;//don't optimize anything out...
}

Answer 2

注意：这个答案不是特定于 python 的，但我认为像 tanh 这样的东西在不同语言之间并没有根本的不同。

Tanh 通常通过定义上限和下限来实现，分别返回 1 和 -1。中间部分用不同的函数近似如下：

 Interval 0  x_small               x_medium               x_large 
  tanh(x) |  x  |  polynomial approx.  |  1-(2/(1+exp(2x)))  |  1

存在精确到单精度浮点和双精度的多项式。 这种算法称为 Cody-Waite 算法。

引用this description（您也可以在那里找到有关数学的更多信息，例如如何确定 x_medium）， Cody and Waite 的有理形式需要 4 次乘法、3 次加法和 1 次单精度除法，以及 7 次乘法、6 次加法和 1 次双精度除法。

对于负 x，你可以计算 |x|并翻转标志。 因此，您需要比较 x 所在的区间，并评估相应的近似值。 总共是：

1. 取x的绝对值
2. 区间的 3 次比较
3. 根据间隔和浮点精度，指数为 0 到几个 FLOPS，请查看this question，了解如何计算指数。
4. 一个比较来决定是否翻转标志。

现在，这是 1993 年的报告，但我认为这里没有太大变化。

Answer 3

这个问题表明它是在机器学习的背景下提出的，因此重点是单精度计算，特别是使用 IEEE-754 binary32 格式。 Asker 表示，NVIDIA GPU 是一个令人感兴趣的平台。我将专注于使用 CUDA 使用这些 GPU，因为我不熟悉 CUDA 的 Python 绑定。

在谈到 FLOPS 时，除了简单的加法和乘法之外，关于如何计算 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 和 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 条浮点指令。

上述代码 cmets 中的错误界限是通过针对所有可能的单精度参数的详尽测试建立的。