HIP math API

HIP math API#

HIP-Clang provides device-callable math operations, supporting most functions available in NVIDIA CUDA.

This section documents:

  • Maximum error bounds for supported HIP math functions

  • Currently unsupported functions

Error bounds on this page are measured in units in the last place (ULPs), representing the absolute difference between a HIP math function result and its corresponding C++ standard library function (e.g., comparing HIP’s sinf with C++’s sinf).

The following C++ example shows a simplified method for computing ULP differences between HIP and standard C++ math functions by first finding where the maximum absolute error occurs.

#include <hip/hip_runtime.h>
#include <iostream>
#include <vector>
#include <cmath>
#include <limits>

#define HIP_CHECK(expression)              \
    {                                      \
        const hipError_t err = expression; \
        if (err != hipSuccess) {           \
            std::cerr << "HIP error: "     \
                      << hipGetErrorString(err) \
                      << " at " << __LINE__ << "\n"; \
            exit(EXIT_FAILURE);            \
        }                                  \
    }

// Simple ULP difference calculator
int64_t ulp_diff(float a, float b) {
    if (a == b) return 0;
    union { float f; int32_t i; } ua{a}, ub{b};

    // For negative values, convert to a positive-based representation
    if (ua.i < 0) ua.i = std::numeric_limits<int32_t>::max() - ua.i;
    if (ub.i < 0) ub.i = std::numeric_limits<int32_t>::max() - ub.i;

    return std::abs((int64_t)ua.i - (int64_t)ub.i);
}

// Test kernel
__global__ void test_sin(float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float x = -M_PI + (2.0f * M_PI * i) / (n - 1);
        out[i] = sin(x);
    }
}

int main() {
    const int n = 1000000;
    const int blocksize = 256;
    std::vector<float> outputs(n);
    float* d_out;

    HIP_CHECK(hipMalloc(&d_out, n * sizeof(float)));
    dim3 threads(blocksize);
    dim3 blocks((n + blocksize - 1) / blocksize);  // Fixed grid calculation
    test_sin<<<blocks, threads>>>(d_out, n);
    HIP_CHECK(hipPeekAtLastError());
    HIP_CHECK(hipMemcpy(outputs.data(), d_out, n * sizeof(float), hipMemcpyDeviceToHost));

    // Step 1: Find the maximum absolute error
    double max_abs_error = 0.0;
    float max_error_output = 0.0;
    float max_error_expected = 0.0;

    for (int i = 0; i < n; i++) {
        float x = -M_PI + (2.0f * M_PI * i) / (n - 1);
        float expected = std::sin(x);
        double abs_error = std::abs(outputs[i] - expected);

        if (abs_error > max_abs_error) {
            max_abs_error = abs_error;
            max_error_output = outputs[i];
            max_error_expected = expected;
        }
    }

    // Step 2: Compute ULP difference based on the max absolute error pair
    int64_t max_ulp = ulp_diff(max_error_output, max_error_expected);

    // Output results
    std::cout << "Max Absolute Error: " << max_abs_error << std::endl;
    std::cout << "Max ULP Difference: " << max_ulp << std::endl;
    std::cout << "Max Error Values -> Got: " << max_error_output
              << ", Expected: " << max_error_expected << std::endl;

    HIP_CHECK(hipFree(d_out));
    return 0;
}

Standard mathematical functions#

The functions in this section prioritize numerical accuracy and correctness, making them well-suited for applications that require high precision and predictable results. Unless explicitly specified, all math functions listed below are available on the device side.

Arithmetic#

Function

Test Range

ULP Difference of Maximum Absolute Error
float abs(float x)
Returns the absolute value of x

x∈[−20,20]

0
float fabsf(float x)
Returns the absolute value of x

x∈[−20,20]

0
float fdimf(float x, float y)
Returns the positive difference between x and y.
x∈[−10,10]
y∈[−3,3]

0
float fmaf(float x, float y, float z)
Returns x⋅y+z as a single operation.
x∈[−100,100]
y∈[−10,10]
z∈[−10,10]

0
float fmaxf(float x, float y)
Determine the maximum numeric value of x and y.
x∈[−10,10]
y∈[−3,3]

0
float fminf(float x, float y)
Determine the minimum numeric value of x and y.
x∈[−10,10]
y∈[−3,3]

0
float fmodf(float x, float y)
Returns the floating-point remainder of x/y.
x∈[−10,10]
y∈[−3,3]

0
float modff(float x, float* iptr)
Break down x into fractional and integral parts.

x∈[−10,10]

0
float remainderf(float x, float y)
Returns single-precision floating-point remainder.
x∈[−10,10]
y∈[−3,3]

0
float remquof(float x, float y, int* quo)
Returns single-precision floating-point remainder and part of quotient.
x∈[−10,10]
y∈[−3,3]

0
float fdividef(float x, float y)
Divide two floating point values.
x∈[−100,100]
y∈[−100,100]

0

Function

Test Range

ULP Difference of Maximum Absolute Error
double abs(double x)
Returns the absolute value of x

x∈[−20,20]

0
double fabs(double x)
Returns the absolute value of x

x∈[−20,20]

0
double fdim(double x, double y)
Returns the positive difference between x and y.
x∈[−10,10]
y∈[−3,3]

0
double fma(double x, double y, double z)
Returns x⋅y+z as a single operation.
x∈[−100,100]
y∈[−10,10]
z∈[−10,10]

0
double fmax(double x, double y)
Determine the maximum numeric value of x and y.
x∈[−10,10]
y∈[−3,3]

0
double fmin(double x, double y)
Determine the minimum numeric value of x and y.
x∈[−10,10]
y∈[−3,3]

0
double fmod(double x, double y)
Returns the floating-point remainder of x/y.
x∈[−10,10]
y∈[−3,3]

0
double modf(double x, double* iptr)
Break down x into fractional and integral parts.

x∈[−10,10]

0
double remainder(double x, double y)
Returns double-precision floating-point remainder.
x∈[−10,10]
y∈[−3,3]

0
double remquo(double x, double y, int* quo)
Returns double-precision floating-point remainder and part of quotient.
x∈[−10,10]
y∈[−3,3]

0

Classification#

Function

Test Range

ULP Difference of Maximum Absolute Error
bool isfinite(float x)
Determine whether x is finite.
x∈[−FLT_MAX,FLT_MAX]
Special values: ±∞, NaN

0
bool isinf(float x)
Determine whether x is infinite.
x∈[−FLT_MAX,FLT_MAX]
Special values: ±∞, NaN

0
bool isnan(float x)
Determine whether x is a NAN.
x∈[−FLT_MAX,FLT_MAX]
Special values: ±∞, NaN

0
bool signbit(float x)
Return the sign bit of x.
x∈[−FLT_MAX,FLT_MAX]
Special values: ±∞, ±0, NaN

0
float nanf(const char* tagp)
Returns “Not a Number” value.
Input strings: "", "1", "2",
"quiet", "signaling", "ind"

0

Function

Test Range

ULP Difference of Maximum Absolute Error
bool isfinite(double x)
Determine whether x is finite.
x∈[−DBL_MAX,DBL_MAX]
Special values: ±∞, NaN

0
bool isin(double x)
Determine whether x is infinite.
x∈[−DBL_MAX,DBL_MAX]
Special values: ±∞, NaN

0
bool isnan(double x)
Determine whether x is a NAN.
x∈[−DBL_MAX,DBL_MAX]
Special values: ±∞, NaN

0
bool signbit(double x)
Return the sign bit of x.
x∈[−DBL_MAX,DBL_MAX]
Special values: ±∞, ±0, NaN

0
double nan(const char* tagp)
Returns “Not a Number” value.
Input strings: "", "1", "2",
"quiet", "signaling", "ind"

0

Error and Gamma#

Function

Test Range

ULP Difference of Maximum Absolute Error
float erff(float x)
Returns the error function of x.

x∈[−4,4]

4
float erfcf(float x)
Returns the complementary error function of x.

x∈[−4,4]

2
float erfcxf(float x)
Returns the scaled complementary error function of x.

x∈[−2,2]

5
float lgammaf(float x)
Returns the natural logarithm of the absolute value of the gamma function of x.

x∈[0.5,20]

4
float tgammaf(float x)
Returns the gamma function of x.

x∈[0.5,15]

6

Function

Test Range

ULP Difference of Maximum Absolute Error
double erf(double x)
Returns the error function of x.

x∈[−4,4]

4
double erfc(double x)
Returns the complementary error function of x.

x∈[−4,4]

2
double erfcx(double x)
Returns the scaled complementary error function of x.

x∈[−2,2]

5
double lgamma(double x)
Returns the natural logarithm of the absolute value of the gamma function of x.

x∈[0.5,20]

2
double tgamma(double x)
Returns the gamma function of x.

x∈[0.5,15]

6

Exponential and Logarithmic#

Function

Test Range

ULP Difference of Maximum Absolute Error
float expf(float x)
Returns ex.

x∈[−10,10]

1
float exp2f(float x)
Returns 2x.

x∈[−10,10]

1
float exp10f(float x)
Returns 10x.

x∈[−4,4]

1
float expm1f(float x)
Returns ln(x−1)

x∈[−10,10]

1
float log10f(float x)
Returns the base 10 logarithm of x.

x∈[10−6,106]

2
float log1pf(float x)
Returns the natural logarithm of x+1.

x∈[−0.9,10]

1
float log2f(float x)
Returns the base 2 logarithm of x.

x∈[10−6,106]

1
float logf(float x)
Returns the natural logarithm of x.

x∈[10−6,106]

2

Function

Test Range

ULP Difference of Maximum Absolute Error
double exp(double x)
Returns ex.

x∈[−10,10]

1
double exp2(double x)
Returns 2x.

x∈[−10,10]

1
double exp10(double x)
Returns 10x.

x∈[−4,4]

1
double expm1(double x)
Returns ln(x−1)

x∈[−10,10]

1
double log10(double x)
Returns the base 10 logarithm of x.

x∈[10−6,106]

1
double log1p(double x)
Returns the natural logarithm of x+1.

x∈[−0.9,10]

1
double log2(double x)
Returns the base 2 logarithm of x.

x∈[10−6,106]

1
double log(double x)
Returns the natural logarithm of x.

x∈[10−6,106]

1

Floating Point Manipulation#

Function

Test Range

ULP Difference of Maximum Absolute Error
float copysignf(float x, float y)
Create value with given magnitude, copying sign of second value.
x∈[−10,10]
y∈[−3,3]

0
float frexpf(float x, int* nptr)
Extract mantissa and exponent of x.

x∈[−10,10]

0
int ilogbf(float x)
Returns the unbiased integer exponent of x.

x∈[0.01,100]

0
float logbf(float x)
Returns the floating point representation of the exponent of x.

x∈[10−6,106]

0
float ldexpf(float x, int exp)
Returns the natural logarithm of the absolute value of the gamma function of x.
x∈[−10,10]
exp∈[−4,4]

0
float nextafterf(float x, float y)
Returns next representable single-precision floating-point value after argument.
x∈[−10,10]
y∈[−3,3]

0
float scalblnf(float x, long int n)
Scale x by 2n.
x∈[−10,10]
n∈[−4,4]

0
float scalbnf(float x, int n)
Scale x by 2n.
x∈[−10,10]
n∈[−4,4]

0

Function

Test Range

ULP Difference of Maximum Absolute Error
double copysign(double x, double y)
Create value with given magnitude, copying sign of second value.
x∈[−10,10]
y∈[−3,3]

0
double frexp(double x, int* nptr)
Extract mantissa and exponent of x.

x∈[−10,10]

0
int ilogb(double x)
Returns the unbiased integer exponent of x.

x∈[0.01,100]

0
double logb(double x)
Returns the floating point representation of the exponent of x.

x∈[10−6,106]

0
double ldexp(double x, int exp)
Returns the natural logarithm of the absolute value of the gamma function of x.
x∈[−10,10]
exp∈[−4,4]

0
double nextafter(double x, double y)
Returns next representable double-precision floating-point value after argument.
x∈[−10,10]
y∈[−3,3]

0
double scalbln(double x, long int n)
Scale x by 2n.
x∈[−10,10]
n∈[−4,4]

0
double scalbn(double x, int n)
Scale x by 2n.
x∈[−10,10]
n∈[−4,4]

0

Hypotenuse and Norm#

Function

Test Range

ULP Difference of Maximum Absolute Error
float hypotf(float x, float y)
Returns the square root of the sum of squares of x and y.
x∈[−10,10]
y∈[0,10]

1
float rhypotf(float x, float y)
Returns one over the square root of the sum of squares of two arguments.
x∈[−100,100]
y∈[−10,100]

1
float norm3df(float x, float y, float z)
Returns the square root of the sum of squares of x, y and z.
All inputs in
[−10,10]

1
float norm4df(float x, float y, float z, float w)
Returns the square root of the sum of squares of x, y, z and w.
All inputs in
[−10,10]

2
float rnorm3df(float x, float y, float z)
Returns one over the square root of the sum of squares of three coordinates of the argument.
All inputs in
[−10,10]

1
float rnorm4df(float x, float y, float z, float w)
Returns one over the square root of the sum of squares of four coordinates of the argument.
All inputs in
[−10,10]

2
float normf(int dim, const float *a)
Returns the square root of the sum of squares of any number of coordinates.
dim∈[2,4]
a[i]∈[−10,10]
Error depends on the number of coordinates
e.g. dim = 2 -> 1
e.g. dim = 3 -> 1
e.g. dim = 4 -> 1
float rnormf(int dim, const float *a)
Returns the reciprocal of square root of the sum of squares of any number of coordinates.
dim∈[2,4]
a[i]∈[−10,10]
Error depends on the number of coordinates
e.g. dim = 2 -> 1
e.g. dim = 3 -> 1
e.g. dim = 4 -> 1

Function

Test Range

ULP Difference of Maximum Absolute Error
double hypot(double x, double y)
Returns the square root of the sum of squares of x and y.
x∈[−10,10]
y∈[0,10]

1
double rhypot(double x, double y)
Returns one over the square root of the sum of squares of two arguments.
x∈[−100,100]
y∈[−10,100]

1
double norm3d(double x, double y, double z)
Returns the square root of the sum of squares of x, y and z.
All inputs in
[−10,10]

1
double norm4d(double x, double y, double z, double w)
Returns the square root of the sum of squares of x, y, z and w.
All inputs in
[−10,10]

2
double rnorm3d(double x, double y, double z)
Returns one over the square root of the sum of squares of three coordinates of the argument.
All inputs in
[−10,10]

1
double rnorm4d(double x, double y, double z, double w)
Returns one over the square root of the sum of squares of four coordinates of the argument.
All inputs in
[−10,10]

1
double norm(int dim, const double *a)
Returns the square root of the sum of squares of any number of coordinates.
dim∈[2,4]
a[i]∈[−10,10]
Error depends on the number of coordinates
e.g. dim = 2 -> 1
e.g. dim = 3 -> 1
e.g. dim = 4 -> 1
double rnorm(int dim, const double *a)
Returns the reciprocal of square root of the sum of squares of any number of coordinates.
dim∈[2,4]
a[i]∈[−10,10]
Error depends on the number of coordinates
e.g. dim = 2 -> 1
e.g. dim = 3 -> 1
e.g. dim = 4 -> 1

Power and Root#

Function

Test Range

ULP Difference of Maximum Absolute Error
float cbrtf(float x)
Returns the cube root of x.

x∈[−100,100]

2
float powf(float x, float y)
Returns xy.
x∈[−4,4]
y∈[−2,2]

1
float powif(float base, int iexp)
Returns the value of first argument to the power of second argument.
base∈[−10,10]
iexp∈[−4,4]

1
float sqrtf(float x)
Returns the square root of x.

x∈[0,100]

1
float rsqrtf(float x)
Returns the reciprocal of the square root of x.

x∈[0.01,100]

1
float rcbrtf(float x)
Returns the reciprocal cube root function.

x∈[−100,100]

1

Function

Test Range

ULP Difference of Maximum Absolute Error
double cbrt(double x)
Returns the cube root of x.

x∈[−100,100]

1
double pow(double x, double y)
Returns xy.
x∈[−4,4]
y∈[−2,2]

1
double powi(double base, int iexp)
Returns the value of first argument to the power of second argument.
base∈[−10,10]
iexp∈[−4,4]

1
double sqrt(double x)
Returns the square root of x.

x∈[0,100]

1
double rsqrt(double x)
Returns the reciprocal of the square root of x.

x∈[0.01,100]

1
double rcbrt(double x)
Returns the reciprocal cube root function.

x∈[−100,100]

1

Rounding#

Function

Test Range

ULP Difference of Maximum Absolute Error
float ceilf(float x)
Returns ceiling of x.

x∈[−4,4]

0
float floorf(float x)
Returns the largest integer less than or equal to x.

x∈[−4,4]

0
long int lroundf(float x)
Round to nearest integer value.

x∈[−4,4]

0
long long int llroundf(float x)
Round to nearest integer value.

x∈[−4,4]

0
long int lrintf(float x)
Round x to nearest integer value.

x∈[−4,4]

0
long long int llrintf(float x)
Round x to nearest integer value.

x∈[−4,4]

0
float nearbyintf(float x)
Round x to the nearest integer.

x∈[−4,4]

0
float roundf(float x)
Round to nearest integer value in floating-point.

x∈[−4,4]

0
float rintf(float x)
Round input to nearest integer value in floating-point.

x∈[−4,4]

0
float truncf(float x)
Truncate x to the integral part.

x∈[−4,4]

0

Function

Test Range

ULP Difference of Maximum Absolute Error
double ceil(double x)
Returns ceiling of x.

x∈[−4,4]

0
double floor(double x)
Returns the largest integer less than or equal to x.

x∈[−4,4]

0
long int lround(double x)
Round to nearest integer value.

x∈[−4,4]

0
long long int llround(double x)
Round to nearest integer value.

x∈[−4,4]

0
long int lrint(double x)
Round x to nearest integer value.

x∈[−4,4]

0
long long int llrint(double x)
Round x to nearest integer value.

x∈[−4,4]

0
double nearbyint(double x)
Round x to the nearest integer.

x∈[−4,4]

0
double round(double x)
Round to nearest integer value in floating-point.

x∈[−4,4]

0
double rint(double x)
Round input to nearest integer value in floating-point.

x∈[−4,4]

0
double trunc(double x)
Truncate x to the integral part.

x∈[−4,4]

0

Trigonometric and Hyperbolic#

Function

Test Range

ULP Difference of Maximum Absolute Error
float acosf(float x)
Returns the arc cosine of x.

x∈[−1,1]

1
float acoshf(float x)
Returns the nonnegative arc hyperbolic cosine of x.

x∈[1,100]

1
float asinf(float x)
Returns the arc sine of x.

x∈[−1,1]

2
float asinhf(float x)
Returns the arc hyperbolic sine of x.

x∈[−10,10]

1
float atanf(float x)
Returns the arc tangent of x.

x∈[−10,10]

2
float atan2f(float x, float y)
Returns the arc tangent of the ratio of x and y.
x∈[−4,4]
y∈[−2,2]

1
float atanhf(float x)
Returns the arc hyperbolic tangent of x.

x∈[−0.9,0.9]

1
float cosf(float x)
Returns the cosine of x.

x∈[−π,π]

1
float coshf(float x)
Returns the hyperbolic cosine of x.

x∈[−5,5]

1
float sinf(float x)
Returns the sine of x.

x∈[−π,π]

1
float sinhf(float x)
Returns the hyperbolic sine of x.

x∈[−5,5]

1
void sincosf(float x, float *sptr, float *cptr)
Returns the sine and cosine of x.

x∈[−3,3]
sin: 1
cos: 1
float tanf(float x)
Returns the tangent of x.

x∈[−1.47π,1.47π]

1
float tanhf(float x)
Returns the hyperbolic tangent of x.

x∈[−5,5]

2
float cospif(float x)
Returns the cosine of π⋅x.

x∈[−0.3,0.3]

1
float sinpif(float x)
Returns the hyperbolic sine of π⋅x.

x∈[−0.625,0.625]

2
void sincospif(float x, float *sptr, float *cptr)
Returns the sine and cosine of π⋅x.

x∈[−0.3,0.3]
sinpi: 2
cospi: 1

Function

Test Range

ULP Difference of Maximum Absolute Error
double acos(double x)
Returns the arc cosine of x.

x∈[−1,1]

1
double acosh(double x)
Returns the nonnegative arc hyperbolic cosine of x.

x∈[1,100]

1
double asin(double x)
Returns the arc sine of x.

x∈[−1,1]

1
double asinh(double x)
Returns the arc hyperbolic sine of x.

x∈[−10,10]

1
double atan(double x)
Returns the arc tangent of x.

x∈[−10,10]

1
double atan2(double x, double y)
Returns the arc tangent of the ratio of x and y.
x∈[−4,4]
y∈[−2,2]

1
double atanh(double x)
Returns the arc hyperbolic tangent of x.

x∈[−0.9,0.9]

1
double cos(double x)
Returns the cosine of x.

x∈[−π,π]

1
double cosh(double x)
Returns the hyperbolic cosine of x.

x∈[−5,5]

1
double sin(double x)
Returns the sine of x.

x∈[−π,π]

1
double sinh(double x)
Returns the hyperbolic sine of x.

x∈[−5,5]

1
void sincos(double x, double *sptr, double *cptr)
Returns the sine and cosine of x.

x∈[−3,3]
sin: 1
cos: 1
double tan(double x)
Returns the tangent of x.

x∈[−1.47π,1.47π]

1
double tanh(double x)
Returns the hyperbolic tangent of x.

x∈[−5,5]

1
double cospi(double x)
Returns the cosine of π⋅x.

x∈[−0.3,0.3]

2
double sinpi(double x)
Returns the hyperbolic sine of π⋅x.

x∈[−0.625,0.625]

2
void sincospi(double x, double *sptr, double *cptr)
Returns the sine and cosine of π⋅x.

x∈[−0.3,0.3]
sinpi: 2
cospi: 2

No C++ STD Implementation#

This table lists HIP device functions that do not have a direct equivalent in the C++ standard library. These functions were excluded from comparison due to the complexity of implementing a precise reference version within the standard library’s constraints.

Function
float j0f(float x)
Returns the value of the Bessel function of the first kind of order 0 for x.
float j1f(float x)
Returns the value of the Bessel function of the first kind of order 1 for x.
float jnf(int n, float x)
Returns the value of the Bessel function of the first kind of order n for x.
float y0f(float x)
Returns the value of the Bessel function of the second kind of order 0 for x.
float y1f(float x)
Returns the value of the Bessel function of the second kind of order 1 for x.
float ynf(int n, float x)
Returns the value of the Bessel function of the second kind of order n for x.
float erfcinvf(float x)
Returns the inverse complementary function of x.
float erfinvf(float x)
Returns the inverse error function of x.
float normcdff(float y)
Returns the standard normal cumulative distribution function.
float normcdfinvf(float y)
Returns the inverse of the standard normal cumulative distribution function.

Function
double j0(double x)
Returns the value of the Bessel function of the first kind of order 0 for x.
double j1(double x)
Returns the value of the Bessel function of the first kind of order 1 for x.
double jn(int n, double x)
Returns the value of the Bessel function of the first kind of order n for x.
double y0(double x)
Returns the value of the Bessel function of the second kind of order 0 for x.
double y1(double x)
Returns the value of the Bessel function of the second kind of order 1 for x.
double yn(int n, double x)
Returns the value of the Bessel function of the second kind of order n for x.
double erfcinv(double x)
Returns the inverse complementary function of x.
double erfinv(double x)
Returns the inverse error function of x.
double normcdf(double y)
Returns the standard normal cumulative distribution function.
double normcdfinv(double y)
Returns the inverse of the standard normal cumulative distribution function.

Unsupported#

This table lists functions that are not supported by HIP.

Function
float cyl_bessel_i0f(float x)
Returns the value of the regular modified cylindrical Bessel function of order 0 for x.
float cyl_bessel_i1f(float x)
Returns the value of the regular modified cylindrical Bessel function of order 1 for x.

Function
double cyl_bessel_i0(double x)
Returns the value of the regular modified cylindrical Bessel function of order 0 for x.
double cyl_bessel_i1(double x)
Returns the value of the regular modified cylindrical Bessel function of order 1 for x.

Intrinsic mathematical functions#

Intrinsic math functions are optimized for performance on HIP-supported hardware. These functions often trade some precision for faster execution, making them ideal for applications where computational efficiency is a priority over strict numerical accuracy. Note that intrinsics are supported on device only.

Floating-point Intrinsics#

Note

Only the nearest-even rounding mode is supported by default on AMD GPUs. The _rz, _ru, and _rd suffixed intrinsic functions exist in the HIP AMD backend if the OCML_BASIC_ROUNDED_OPERATIONS macro is defined.

Single precision intrinsics mathematical functions#

Function

Test Range

ULP Difference of Maximum Absolute Error
float __cosf(float x)
Returns the fast approximate cosine of x.

x∈[−π,π]

4
float __exp10f(float x)
Returns the fast approximate for 10 x.

x∈[−4,4]

18
float __expf(float x)
Returns the fast approximate for e x.

x∈[−10,10]

6
float __fadd_rn(float x, float y)
Add two floating-point values in round-to-nearest-even mode.
x∈[−1000,1000]
y∈[−1000,1000]

0
float __fdiv_rn(float x, float y)
Divide two floating-point values in round-to-nearest-even mode.
x∈[−100,100]
y∈[−100,100]

0
float __fmaf_rn(float x, float y, float z)
Returns x × y + z as a single operation in round-to-nearest-even mode.
x∈[−100,100]
y∈[−10,10]
z∈[−10,10]

0
float __fmul_rn(float x, float y)
Multiply two floating-point values in round-to-nearest-even mode.
x∈[−100,100]
y∈[−100,100]

0
float __frcp_rn(float x, float y)
Returns 1 / x in round-to-nearest-even mode.

x∈[−100,100]

0
float __frsqrt_rn(float x)
Returns 1 / √x in round-to-nearest-even mode.

x∈[0.01,100]

1
float __fsqrt_rn(float x)
Returns √x in round-to-nearest-even mode.

x∈[0,100]

1
float __fsub_rn(float x, float y)
Subtract two floating-point values in round-to-nearest-even mode.
x∈[−1000,1000]
y∈[−1000,1000]

0
float __log10f(float x)
Returns the fast approximate for base 10 logarithm of x.

x∈[10−6,106]

2
float __log2f(float x)
Returns the fast approximate for base 2 logarithm of x.

x∈[10−6,106]

1
float __logf(float x)
Returns the fast approximate for natural logarithm of x.

x∈[10−6,106]

2
float __powf(float x, float y)
Returns the fast approximate of x y.
x∈[−4,4]
y∈[−2,2]

1
float __saturatef(float x)
Clamp x to [+0.0, 1.0].

x∈[−2,3]

0
float __sincosf(float x, float* sinptr, float* cosptr)
Returns the fast approximate of sine and cosine of x.

x∈[−3,3]
sin: 18
cos: 4
float __sinf(float x)
Returns the fast approximate sine of x.

x∈[−π,π]

18
float __tanf(float x)
Returns the fast approximate tangent of x.

x∈[−1.47π,1.47π]

1
Double precision intrinsics mathematical functions#

Function

Test Range

ULP Difference of Maximum Absolute Error
double __dadd_rn(double x, double y)
Add two floating-point values in round-to-nearest-even mode.
x∈[−1000,1000]
y∈[−1000,1000]

0
double __ddiv_rn(double x, double y)
Divide two floating-point values in round-to-nearest-even mode.
x∈[−100,100]
y∈[−100,100]

0
double __dmul_rn(double x, double y)
Multiply two floating-point values in round-to-nearest-even mode.
x∈[−100,100]
y∈[−100,100]

0
double __drcp_rn(double x, double y)
Returns 1 / x in round-to-nearest-even mode.

x∈[−100,100]

0
double __dsqrt_rn(double x)
Returns √x in round-to-nearest-even mode.

x∈[0,100]

0
double __dsub_rn(double x, double y)
Subtract two floating-point values in round-to-nearest-even mode.
x∈[−1000,1000]
y∈[−1000,1000]

0
double __fma_rn(double x, double y, double z)
Returns x × y + z as a single operation in round-to-nearest-even mode.
x∈[−100,100]
y∈[−10,10]
z∈[−10,10]

0

Integer intrinsics#

This section covers HIP integer intrinsic functions. ULP error values are omitted since they only apply to floating-point operations, not integer arithmetic.

Integer intrinsics mathematical functions#

Function
unsigned int __brev(unsigned int x)
Reverse the bit order of a 32 bit unsigned integer.
unsigned long long int __brevll(unsigned long long int x)
Reverse the bit order of a 64 bit unsigned integer.
unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int z)
Return selected bytes from two 32-bit unsigned integers.
unsigned int __clz(int x)
Return the number of consecutive high-order zero bits in 32 bit integer.
unsigned int __clzll(long long int x)
Return the number of consecutive high-order zero bits in 64 bit integer.
unsigned int __ffs(int x) [1]
Returns the position of the first set bit in a 32 bit integer.
Note: if x is 0, will return 0
unsigned int __ffsll(long long int x) [1]
Returns the position of the first set bit in a 64 bit signed integer.
Note: if x is 0, will return 0
unsigned int __fns32(unsigned long long mask, unsigned int base, int offset)
Find the position of the n-th set to 1 bit in a 32-bit integer.
Note: this intrinsic is emulated via software, so performance can be potentially slower
unsigned int __fns64(unsigned long long int mask, unsigned int base, int offset)
Find the position of the n-th set to 1 bit in a 64-bit integer.
Note: this intrinsic is emulated via software, so performance can be potentially slower
unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi and lo, shift left by shift & 31 bits, return the most significant 32 bits.
unsigned int __funnelshift_lc(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi and lo, shift left by min(shift, 32) bits, return the most significant 32 bits.
unsigned int __funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi and lo, shift right by shift & 31 bits, return the least significant 32 bits.
unsigned int __funnelshift_rc(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi and lo, shift right by min(shift, 32) bits, return the least significant 32 bits.
unsigned int __hadd(int x, int y)
Compute average of signed input arguments, avoiding overflow in the intermediate sum.
unsigned int __rhadd(int x, int y)
Compute rounded average of signed input arguments, avoiding overflow in the intermediate sum.
unsigned int __uhadd(int x, int y)
Compute average of unsigned input arguments, avoiding overflow in the intermediate sum.
unsigned int __urhadd (unsigned int x, unsigned int y)
Compute rounded average of unsigned input arguments, avoiding overflow in the intermediate sum.
int __sad(int x, int y, int z)
Returns |x−y|+z, the sum of absolute difference.
unsigned int __usad(unsigned int x, unsigned int y, unsigned int z)
Returns |x−y|+z, the sum of absolute difference.
unsigned int __popc(unsigned int x)
Count the number of bits that are set to 1 in a 32 bit integer.
unsigned int __popcll(unsigned long long int x)
Count the number of bits that are set to 1 in a 64 bit integer.
int __mul24(int x, int y)
Multiply two 24bit integers.
unsigned int __umul24(unsigned int x, unsigned int y)
Multiply two 24bit unsigned integers.
int __mulhi(int x, int y)
Returns the most significant 32 bits of the product of the two 32-bit integers.
unsigned int __umulhi(unsigned int x, unsigned int y)
Returns the most significant 32 bits of the product of the two 32-bit unsigned integers.
long long int __mul64hi(long long int x, long long int y)
Returns the most significant 64 bits of the product of the two 64-bit integers.
unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y)
Returns the most significant 64 bits of the product of the two 64 unsigned bit integers.