Skip to main content
Fix typo
Source Link
njuffa
  • 540
  • 3
  • 12

The same accuracy and behavior can be achieved by using a simple mathematical transformation that eliminates $\mathrm{tahn}$$\mathrm{tanh}$, and considering that GPUs typically offer a fused multiply-add (FMA) as well as a fast reciprocal, which one would want to utilize. Exemplary CUDA code looks as follows:

The same accuracy and behavior can be achieved by using a simple mathematical transformation that eliminates $\mathrm{tahn}$, and considering that GPUs typically offer a fused multiply-add (FMA) as well as a fast reciprocal, which one would want to utilize. Exemplary CUDA code looks as follows:

The same accuracy and behavior can be achieved by using a simple mathematical transformation that eliminates $\mathrm{tanh}$, and considering that GPUs typically offer a fused multiply-add (FMA) as well as a fast reciprocal, which one would want to utilize. Exemplary CUDA code looks as follows:

Add note on machine-specific performance enahncement.
Source Link
njuffa
  • 540
  • 3
  • 12

As a machine-specific performance enhancement, expf() could be replaced by the device intrinsic __expf().

As a machine-specific performance enhancement, expf() could be replaced by the device intrinsic __expf().

Add third solution with reduced accuracy (but smooth tail) and improved performance
Source Link
njuffa
  • 540
  • 3
  • 12

Unfortunately, this accurate solution is achieved at the cost of a significant drop in performance. If one is willing to accept reduced accuracy while still achieving a smoothly decaying left tail, the following interpolation scheme, again based on $f(x) \approx x\exp(x)$, recovers much of the performance:

__device__ float my_mishf (float x)
{
    float r;
    float e = expf (x);
    if (x >= -6.0625f) {
        r = 1.0f / fmaf (fmaf (-0.5f, e, -1.0f), e, -1.0f);
        r = fmaf (r, x, x);
    } else {
        r = fmaf (-0.5f, e, 1.0f);
        r = r * x * e;
    }
    return r;
}

Unfortunately, this accurate solution is achieved at the cost of a significant drop in performance. If one is willing to accept reduced accuracy while still achieving a smoothly decaying left tail, the following interpolation scheme, again based on $f(x) \approx x\exp(x)$, recovers much of the performance:

__device__ float my_mishf (float x)
{
    float r;
    float e = expf (x);
    if (x >= -6.0625f) {
        r = 1.0f / fmaf (fmaf (-0.5f, e, -1.0f), e, -1.0f);
        r = fmaf (r, x, x);
    } else {
        r = fmaf (-0.5f, e, 1.0f);
        r = r * x * e;
    }
    return r;
}
Fix errors in comments; minor word smithing
Source Link
njuffa
  • 540
  • 3
  • 12
Loading
Fix error in description. Use more accurate rational approximation.
Source Link
njuffa
  • 540
  • 3
  • 12
Loading
Minor fixes.
Source Link
njuffa
  • 540
  • 3
  • 12
Loading
Add link to OP's reference implementation
Source Link
njuffa
  • 540
  • 3
  • 12
Loading
Source Link
njuffa
  • 540
  • 3
  • 12
Loading