nvcc Kokkos_ENABLE_DEBUG_BOUNDS_CHECK compile perf fix for Kokkos_Cor…

…eUnitTest_Default

core/src/Cuda/Kokkos_Cuda_abort.hpp

@@ -28,7 +28,7 @@ extern "C" {
 /*  Cuda runtime function, declared in <crt/device_runtime.h>
  *  Requires capability 2.x or better.
  */
-extern __device__ void __assertfail(const void *message, const void *file,
+__device__ [[noreturn]] void __assertfail(const void *message, const void *file,
                                     unsigned int line, const void *function,
                                     size_t charsize);
 }
@@ -38,25 +38,11 @@ namespace Impl {
 
 // required to workaround failures in random number generator unit tests with
 // pre-volta architectures
-#if defined(KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK)
-__device__ inline void cuda_abort(const char *const message) {
-#else
-[[noreturn]] __device__ inline void cuda_abort(const char *const message) {
-#endif
+__device__ __noinline__ [[noreturn]] static  void cuda_abort(const char *const message) {
   const char empty[] = "";
 
   __assertfail((const void *)message, (const void *)empty, (unsigned int)0,
                (const void *)empty, sizeof(char));
-
-  // This loop is never executed. It's intended to suppress warnings that the
-  // function returns, even though it does not. This is necessary because
-  // __assertfail is not marked as [[noreturn]], even though it does not return.
-  //  Disable with KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK to workaround failures
-  //  in random number generator unit tests with pre-volta architectures
-#if !defined(KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK)
-  while (true)
-    ;
-#endif
 }
 
 }  // namespace Impl