Kokkos::ArithTraits<double>::nan() is very slow

[kokkos]

Kokkos::ArithTraits::nan() is very slow

from kokkos-kernels.

[crtrott]

Just try out what is fastest and do that. Though I hope its the std::numeric_limits solution ;-)

[mhoemmen]

Kokkos::ArithTraits needs to work on the GPU, so you can't use std::numeric_limits if __CUDA_ARCH__ is defined. It could be that this works now, given how CUDA's constexpr support has been evolving. However, you'll have to try it to be sure.

If std::numeric_limits doesn't work, you could try this:
https://stackoverflow.com/questions/15514286/way-to-get-floating-point-special-values-in-cuda

[ibaned]

Actually, this is the code right now, so the CUDA part is already separate:

#ifdef __CUDA_ARCH__
    return CUDART_NAN;
    //return nan (); // this returns 0 ???
#else
    // http://pubs.opengroup.org/onlinepubs/009696899/functions/nan.html
    return strtod ("NAN", (char**) NULL);
#endif // __CUDA_ARCH__

So we only have to worry about the slow host code, which has no GPU restrictions.

[hcedwar]

Suggest the strtod only be called once to initialize a singleton:

#ifdef __CUDA_ARCH__
__device__ inline
double kokkos_kernels_nan_function() // whatever 
{  return CUDART_NAN ; }
#else
inline
double kokkos_kernels_nan_function() // whatever 
{ static double x = strtod("NAN",(char**)NULL); return x ; }
#endif

[ibaned]

This is quite similar to what Teuchos does. They also use a singleton, just a different initialization.

Any reason not to use the numeric_limits quiet_NAN to initialize the singleton instead of strtod ?

[ibaned]

Actually, std::numeric_limits<double>::quiet_NaN() is constexpr in C++11, so calling it directly might be faster than loading a singleton in that case.

[ibaned]

Could someone add me to this repository so I can assign this to myself?

[crtrott]

Ah this is Kokkos-Kernels totally overlooked this.

[mhoemmen]

std::numeric_limits should work. We could even make ArithTraits::nan constexpr in the non-CUDA case (and eventually, even in the CUDA case, since CUDA looks like it's starting to get more constexpr support).

Singletons could be bad if calling this in a thread-parallel context; we would want pthread_once in that case.

[ibaned]

trilinos/Trilinos#1490 fixes this in Trilinos, including the addition of a performance test in the affected MiniTensor package to measure the difference.

#36 fixes this in the KokkosKernels develop branch.

They are separate since KokkosKernels is unlikely to promote into Trilinos in the near term, and MiniTensor needs this now.

[mhoemmen]

I submitted a trilinos/Trilinos#1490 review; thanks!