CODEWITHSUNDEEP
cudaintrinsics
Fabio T.
Fabio T.

Reputation: 129

Semantics of __ddiv_ru

From the documentation of __ddiv_ru I expect that the following code result is ceil(8/32) = 1.0, instead I obtain 0.25.

#include <iostream>

using namespace std;

__managed__ double x;
__managed__ double y;
__managed__ double r;

__global__ void ceilDiv()
{
    r = __ddiv_ru(x,y);
}

int  main()
{
    x = 8;
    y = 32;
    r = -1;

    ceilDiv<<<1,1>>>();
    cudaDeviceSynchronize();

    cout << "The ceil of " << x << "/" << y << " is " << r << endl;

    return 1;
}

What am I missing?

Upvotes: 0

Views: 149

Answers (1)

talonmies
talonmies

Reputation: 72349

The result you are obtaining is correct.

The intrinsic you are using implements double precision division with a specific IEEE 754-2008 rounding mode for the unit in the last place (ULP) of the significand. This controls what happens when a result cannot be exactly represented in the selected format. In this case you have selected round up, which means the last digit of the significand produced in the division result is rounded up (toward +∞). In your case all rounding modes should produce the same result because the result can be exactly represented in IEEE 754 binary64 format (it is a round power of 2).

Please read everything here before writing any more floating point code.

Upvotes: 4

Related Questions