仅对数底数为 2 时的自然对数实现

Posted 2023-02-16

【中文标题】仅对数底数为 2 时的自然对数实现

【英文标题】：Natural logarithm implementation when only logarithm base 2 is available

【发布时间】：2012-11-16 16:45:22

【问题描述】：

我正在尝试使用 PTX 实现自然对数​​。 PTX 本身仅提供lg2.approx.f32，它实现了以 2 为底的对数。因此，应用简单的数学运算，只需将以 2 为底的对数乘以欧拉数 e 的以 2 为底的对数即可得到自然对数：

log_e(a) = log_2(a) / lg_2(e)

1/lg_2(e) 的第一个近似值是0.693147。所以，我会乘以这个数字。

我让nvcc 将log 函数（来自CUDA C）编译成PTX（请在下面找到输出）。我可以看到最后与数字常数相乘。但是还有很多事情要做。这很重要吗？有人能解释一下为什么会有这么多开销吗？

    .entry _Z6kernelPfS_ (
        .param .u64 __cudaparm__Z6kernelPfS__out,
        .param .u64 __cudaparm__Z6kernelPfS__in)
    
    .reg .u32 %r<13>;
    .reg .u64 %rd<4>;
    .reg .f32 %f<48>;
    .reg .pred %p<4>;
    .loc    14  3   0
$LDWbegin__Z6kernelPfS_:
    .loc    14  5   0
    ld.param.u64    %rd1, [__cudaparm__Z6kernelPfS__in];
    ld.global.f32   %f1, [%rd1+0];
    .loc    16  9365    0
    mov.f32     %f2, 0f00000000;        // 0
    set.gt.u32.f32  %r1, %f1, %f2;
    neg.s32     %r2, %r1;
    mov.f32     %f3, 0f7f800000;        // ((1.0F)/(0.0F))
    set.lt.u32.f32  %r3, %f1, %f3;
    neg.s32     %r4, %r3;
    and.b32     %r5, %r2, %r4;
    mov.u32     %r6, 0;
    setp.eq.s32     %p1, %r5, %r6;
    @%p1 bra    $Lt_0_2306;
    .loc    16  8512    0
    mov.b32     %r7, %f1;
    and.b32     %r8, %r7, -2139095041;
    or.b32  %r9, %r8, 1065353216;
    mov.b32     %f4, %r9;
    mov.f32     %f5, %f4;
    .loc    16  8513    0
    shr.u32     %r10, %r7, 23;
    sub.u32     %r11, %r10, 127;
    mov.f32     %f6, 0f3fb504f3;        // 1.41421
    setp.gt.f32     %p2, %f4, %f6;
    @!%p2 bra   $Lt_0_2562;
    .loc    16  8515    0
    mov.f32     %f7, 0f3f000000;        // 0.5
    mul.f32     %f5, %f4, %f7;
    .loc    16  8516    0
    add.s32     %r11, %r11, 1;
$Lt_0_2562:
    .loc    16  8429    0
    mov.f32     %f8, 0fbf800000;        // -1
    add.f32     %f9, %f5, %f8;
    mov.f32     %f10, 0f3f800000;       // 1
    add.f32     %f11, %f5, %f10;
    neg.f32     %f12, %f9;
    div.approx.f32  %f13, %f9, %f11;
    mul.rn.f32  %f14, %f12, %f13;
    add.rn.f32  %f15, %f9, %f14;
    mul.f32     %f16, %f15, %f15;
    mov.f32     %f17, 0f3b2063c3;       // 0.00244735
    mov.f32     %f18, %f17;
    mov.f32     %f19, %f16;
    mov.f32     %f20, 0f3c4c4be0;       // 0.0124693
    mov.f32     %f21, %f20;
    mad.f32 %f22, %f18, %f19, %f21;
    mov.f32     %f23, %f22;
    mov.f32     %f24, %f23;
    mov.f32     %f25, %f16;
    mov.f32     %f26, 0f3daaab50;       // 0.0833346
    mov.f32     %f27, %f26;
    mad.f32 %f28, %f24, %f25, %f27;
    mov.f32     %f29, %f28;
    mul.f32     %f30, %f16, %f29;
    mov.f32     %f31, %f30;
    mov.f32     %f32, %f15;
    mov.f32     %f33, %f14;
    mad.f32 %f34, %f31, %f32, %f33;
    mov.f32     %f35, %f34;
    cvt.rn.f32.s32  %f36, %r11;
    mov.f32     %f37, %f36;
    mov.f32     %f38, 0f3f317218;       // 0.693147
    mov.f32     %f39, %f38;
    add.f32     %f40, %f9, %f35;
    mov.f32     %f41, %f40;
    mad.f32 %f42, %f37, %f39, %f41;
    mov.f32     %f43, %f42;
    .loc    16  8523    0
    mov.f32     %f44, %f43;
    bra.uni     $Lt_0_2050;
$Lt_0_2306:
    .loc    16  8526    0
    lg2.approx.f32  %f45, %f1;
    mov.f32     %f46, 0f3f317218;       // 0.693147
    mul.f32     %f44, %f45, %f46;
$Lt_0_2050:
    .loc    14  5   0
    ld.param.u64    %rd2, [__cudaparm__Z6kernelPfS__out];
    st.global.f32   [%rd2+0], %f44;
    .loc    14  6   0
    exit;
$LDWend__Z6kernelPfS_:
     // _Z6kernelPfS_

* 编辑 *

只是为了完整。这里是我编译成上述 PTX 的 CUDA C 内核：

__global__ void kernel(float *out, float *in)

  *out = log( *in );

【问题讨论】：

您可以在 math_functions.h 中查看 logf() 的 C 源代码。 GPU 硬件内置的 log2 相对误差较大，接近 1.0，因此使用纯软件逼近代替硬件，以在整个函数域上实现较小的 ulp 误差。如果您只想要硬件 log2 加上反乘，请使用 __logf() 而不是 logf()。

好主意。我刚试过logf。不幸的是，PTX 没有变化。似乎他们无论如何都使用近似值。您是否看到在什么情况下代码分支到“硬件 lg2 和反向乘法”。对我来说它并不全面。

__logf() 带有两个前导下划线是直接映射到硬件近似值的“快速数学”版本。请参阅文件 device_functions.h 中的源代码。常规 logf() 函数仅使用 __logf() 来处理异常情况，这可以通过查看 math_functions.h 中的源代码轻松看出。这两个文件在任何安装了 CUDA 的机器上都很容易获得，因为此时 CUDA 标准数学库只是头文件的集合。

是的，__logf() 只产生它的硬件 + 乘法版本。那是标签Lt_0_2306之后的代码。回到我的问题：该标签的近似代码中有一个分支命令。如果 Pade 近似失效，你会说这个分支被采用吗？作为故障保护版本。

来源： if ((a > CUDART_ZERO_F) && (a

【参考方案1】：

该函数似乎通过将指数部分设置为 1（有效缩放到范围 1<x<2）然后计算 3 次多项式近似值来计算浮点数的 log2。 编辑：似乎是理性的Pade approximation，因为 log(1+x) 的泰勒级数收敛很差。从而计算倒数。

可能最多只能删除少量指令。 （代码乘以 0.5，而不是从指数中减去这些琐碎的东西。例如，测试参数 x

【讨论】：

NVIDIA GPU 上的单精度浮点计算通常不比整数计算贵。代码效率分析必须在 SASS 级别进行，这是在硬件上执行的。 PTX 只是可移植的中间代码，并没有完全优化。由于这个问题促使我查看 logf() 的源代码，我借此机会进一步简化了代码:-)

阿基，非常感谢！几个月来我一直在寻找这个。