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

[ibaned]

I just got a report from Albany users (@lxmota and @calleman21). They switched from Teuchos::ScalarTraits<double>::nan() to Kokkos::ArithTraits<double>::nan() which slowed down their entire application by over 3X because all variables are initialized to NaN using this function.

KokkosKernels uses strtod() to implement this function (on the host), while Teuchos returns a global variable which is initialized to 0.0/0.0. @lxmota also recommended that we simply call std::numeric_traits<double>::quiet_NaN().

All the above also applies to float.

I think we should switch to either what Teuchos does or the quiet_NaN() from the standard library.

@crtrott @mhoemmen any thoughts?

I'll pick one of these and submit a PR soon.

[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.

[crtrott]

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

[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!