tanh 需要多少个 FLOP？

Question

我想计算 LeNet-5 的每一层有多少次失败（paper) needs. Some papers give FLOPs for other architectures in total (1, 2, 3) However, those papers don't give details on how to compute the number of FLOPs and I have no idea how many FLOPs are necessary for the non-linear activation functions. For example, how many FLOPs are necessary to calculate tanh(x)？

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

Answer 1

注意：这个答案不是 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 和 Waite 的有理式要求单精度四乘三加一除，双精度七乘六加一除

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

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

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

Answer 2

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

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

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

我在 Intel-processor 上的（草率）实验产生了以下结果，这给出了一个粗略的想法：

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

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

相信有一个叫做F2XM1 which computes 2^x-1 for -1.0<x<1.0, which could be used for computing tanh, at least for some range. However if agner's tables的Intel-instruction，这个操作的成本大约是60 FLOPs。

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

如您在表中所见，每次操作的最大成本约为 10 个时钟（float-operation 的成本约为 1 个时钟）。

我猜你有机会通过使用 -ffast-math 编译器选项赢得一些 FLOPs，这会导致更快但不太精确的程序（这是 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 3

这个问题表明它是在机器学习的背景下提出的，因此重点是单精度计算，特别是使用 IEEE-754 binary32 格式。 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;
}

在计算能力 < 7.5 的旧 GPU 架构上，我们可以通过代数变换和使用机器指令 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 目标生成的机器代码表明该分支已保留。这意味着一般而言，在所有活动线程中，一些线程将遵循分支的一侧，而其余线程将遵循分支的另一侧，需要后续同步。因此，为了对所需的计算工作量有一个现实的想法，我们需要结合 both 执行路径的浮点指令计数。这得出十三个浮点指令。

上面代码注释中的错误界限是通过针对所有可能的单精度参数的详尽测试确定的。

tanh 需要多少个 FLOP？

How many FLOPs does tanh need?

python

flops

tensorflow