Assigning NaN to a variable in a device function

I am looking to assign NaN to a variable in a device function in my CUDA code

Example code: if isnan(A[i]) { B= NaN;}

I attempted the above, but I get an error message saying NaN not defined.
I tried changing it to Nan and then nan. Didn’t work for either. Same error message as above.

The C++ standard defines two ways of generating a NaN that it inherits from C (for C++ specific functionality, see the answer by @Robert_Crovella):

  1. A predefined macro NAN, which has type float, and evaluates to an implementation-specific generic NaN. This may not be the same encoding resulting from INFINITY / INFINITY, for example: On a x86-64 Linux system with gcc, NAN is likely encoded as 0x7fc00000, while an invalid operation like INFINITY/ INFINITY returns a QNaN encoding 0xffc00000 (called INDEFINITE in Intel parlance).

  2. A math function nan(), with single-precision variant nanf(), that takes a string argument that translates to a quiet NaN encoding in an implementation-defined manner. Passing an empty string results in a generic NaN being returned.

The GPU hardware does not support NaN “payloads” for float operands, so nanf("") is a cheap operation that simply returns the canonical single-precision quiet NaN: __int_as_float (0x7fffffff). For double precision, NaN payloads are supported, and nan() uses a small parser to convert the passed string into a floating-point operand. This is reasonably expensive. You could experiment with things like nan ("123") or nan ("0xfe")
or nan ("071") and examine what they return.

#include <cstdio>
#include <cstdlib>
#include <cmath>

#define VARIANT  (0)

__global__ void kernel (float x, float *r)
{
#if VARIANT == 0
    *r = isnan (x) ? NAN : x;
#elif VARIANT == 1
    *r = isnan (x) ? nanf("") : x;
#else
#error unknown VARIANT
#endif    
}

int main (void)
{
    float arg, res, *res_d = 0;
    cudaMalloc ((void**)&res_d, sizeof (*res_d));
    arg = -5.0f;
    kernel<<<1,1>>>(arg, res_d);
    cudaMemcpy (&res, res_d, sizeof res, cudaMemcpyDeviceToHost);
    printf ("res = %f\n", res);
    arg = sqrtf (arg);
    kernel<<<1,1>>>(arg, res_d);
    cudaMemcpy (&res, res_d, sizeof res, cudaMemcpyDeviceToHost);
    printf ("res = %f\n", res);
    return EXIT_SUCCESS;
}

The libcu++ library provides device NaN functionality similar to what std::numeric_limits provides for host code.