Testing for infinity in CUDA

Question

In a CUDA program, I recently switched from testing for inifinity using

return x==INFINITY || x==-INFINITY;

where INFINITY is from math.h, to

return !isfinite(x);

and was quite surprised to get different results. gnu.org suggests that they actually should behave similarly. Am I missing something? Is it not allowed to use INFINITY in a CUDA kernel?

Edit: I just discovered isinf and noticed that checking using

return isinf(x);

gives the same result as the INFINITY check. Why isn't isfinite(x)==!isinf(x)?

For what values did the results differ? What were the results? What were the expected results? — James McNellis, Dec 03 '10 at 22:08
i'm not sure yet, but it seems isfinite is more strict than the INFINITY test. — hannes, Dec 03 '10 at 22:14
@stephen canon: granted, but the Cg docs don't mention the INFINITY macro from math.h, the difference to which I am asking about. Apparently the INFINITY macro does work like isinf (see my edit), so it has some relevance. — hannes, Dec 03 '10 at 22:33

score 9 · Accepted Answer · answered Dec 03 '10 at 22:27

9

isfinite(a) is the same as !isnan(a) && !isinf(a). If x is NaN, then both isfinite(x) and isinf(x) are false.

answered Dec 03 '10 at 22:27

Stephen Canon

103,815
19
183
269

My check for NaN `(x!=x)` passes, so according to this test there are no NaNs. However, if Crashworks is right, then this might actually be the solution. – hannes Dec 03 '10 at 22:38
@hannes: what GPU and version of CUDA are you using? – Stephen Canon Dec 03 '10 at 22:55
3

@hannes: Make sure `(x!=x)` isn't being optimized away by your compiler. – Billy ONeal Dec 03 '10 at 22:57
@Billy ONeal: the nan-test works correctly when inject NaNs manually. – hannes Dec 03 '10 at 23:06

score 3 · Answer 2 · answered Dec 03 '10 at 22:51

3

isinf() checks only for +INFINITY or -INFINITY.

!isfinite() checks for +INFINITY, -INFINITY or NaN.

answered Dec 03 '10 at 22:51

caf

233,326
40
323
462

score 1 · Answer 3 · answered Dec 03 '10 at 22:08

1

Floating-point comparisons aren't necessarily valid. For example, it's possible that (1.0f + 3.0f != 2.0f + 2.0f).It's perfectly possible that isfinite considers values less than a specific constant apart to be equal to INFINITE or -INFINITE, whereas you wrote a literal equality.

answered Dec 03 '10 at 22:08

Puppy

144,682
38
256
465

11

It is **not** possible that `1.0f + 3.0f != 2.0f + 2.0f` on any system with the IEEE-754 arithmetic (or even on systems that aren't-quite-IEEE-754-compliant like some GPUs). Floating-point is not black magic. There are well-defined rules for when it rounds and when it does not. – Stephen Canon Dec 03 '10 at 22:21

Crashworks · Answer 4 · 2010-12-07T13:05:25.387

1

Many GPUs and SIMD units are not totally IEEE754 compliant, especially for edge cases around infinities and NaNs. Just last night I noticed that a particular vector processor I worked with claimed that ∞+1 ≠ ∞ , and x == x even for x &in; NaN.

edited Dec 07 '10 at 13:05

answered Dec 03 '10 at 22:12

Crashworks

40,496
12
101
170

score 0 · Answer 5 · edited May 23 '17 at 12:34

In the recent post Checking if a matrix contains nans or infinite values in CUDA, Robert Crovella has suggested using isinf() to check for infinite values in CUDA.

Below I'm providing an example of checking for infinite values in an array using isinf() and by exploiting CUDA Thrust. Perhaps it could be useful as reference for other users. The example below is equivalent to Matlab's d_result=isinf(d_data);. It is different from the example I posted for the question quoted above in that the present one checks for each individual element is infinite while the other one checked if the whole array contained at least one NaN and was equivalent to Matlab's sum(isnan(d_data));.

#include <thrust/sequence.h>

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust\device_vector.h>
#include <thrust\reduce.h>

#include <float.h>

// --- Operator for testing inf values
struct isinf_test { 
    __host__ __device__ bool operator()(const float a) const {
        return isinf(a);
    }
};

void main(){

    const int N = 10;

    thrust::host_vector<float> h_data(N);
    for (int i=0; i<N; i++)
        h_data[i] = rand()/RAND_MAX;

    h_data[0] = FLT_MAX/FLT_MIN;

    thrust::device_vector<float> d_data(h_data);
    thrust::device_vector<float> d_result(h_data);

    thrust::transform(d_data.begin(), d_data.end(), d_result.begin(), isinf_test());

    for (int i=0; i<N; i++) {
        float val = d_result[i];
        printf("Isinf test for element number %i equal to %f\n",i,val);
    }

    getchar();

}

Testing for infinity in CUDA

5 Answers5