11

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

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

4

3 回答 3

10

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

  1. 对于 x大于 22.0 和双精度的值,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).

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

我在英特尔处理器上的(草率)实验产生了以下结果,这给出了一个粗略的想法:

在此处输入图像描述

所以对于非常小的数字和大于 22 的数字,几乎没有成本,对于最多0.1我们支付 6 FLOPS 的数字,成本上升到每次计算大约 20 FLOPS tanh

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


有一个 Intel 指令称为F2XM1which computes 2^x-1for -1.0<x<1.0,它可以用于计算tanh,至少在某个范围内。但是,如果要相信agner 的表格,则此操作的成本约为 60 FLOP。


另一个问题是向量化——据我所知,正常的 glibc 实现没有向量化。因此,如果您的程序使用矢量化并且必须使用未矢量化的tanh实现,它将更加减慢程序的速度。为此,intel 编译器具有 mkl-library,其中向量化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...
}
于 2017-03-25T20:51:25.103 回答
8

注意:这个答案不是特定于 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 算法。

引用这个描述(您也可以在那里找到有关数学的更多信息,例如如何确定 x_medium),Cody 和 Waite 的有理形式需要四次乘法、三次加法和一次单精度除法,以及七次乘法、六次加法和双精度除法。

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

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

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

于 2017-03-25T10:32:41.570 回答
1

这个问题表明它是在机器学习的背景下提出的,因此重点是单精度计算,特别是使用 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 条浮点指令。

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

于 2021-04-07T20:00:27.723 回答