(Rebase) Partial fix to compile time issues w/nvcc + Kokkos_ENABLE_DEBUG_BOUNDS_CHECK

[nmm0]

Rebase of #6596

For reference when configured with:

  CMAKE_BUILD_TYPE="Debug"
  CMAKE_COMPILE_WARNING_AS_ERROR="ON"
  CMAKE_CXX_COMPILER="kokkos/bin/nvcc_wrapper"
  CMAKE_CXX_STANDARD="17"
  CMAKE_EXE_LINKER_FLAGS="-rdynamic"
  CMAKE_EXPORT_COMPILE_COMMANDS="ON"
  CMAKE_INSTALL_PREFIX:PATH="..."
  CMAKE_POSITION_INDEPENDENT_CODE="ON"
  CMAKE_VERBOSE_MAKEFILE:BOOL="OFF"
  CUDAToolkit_BIN_DIR=".../spack-env/view/gcc-12.2.0/cuda/11.6.0/bin"
  CUDAToolkit_ROOT=".../spack-env/view/gcc-12.2.0/cuda/11.6.0"
  Kokkos_ENABLE_CUDA="ON"
  Kokkos_ENABLE_CUDA_CONSTEXPR="OFF"
  Kokkos_ENABLE_CUDA_LAMBDA="ON"
  Kokkos_ENABLE_DEBUG="ON"
  Kokkos_ENABLE_DEBUG_BOUNDS_CHECK="ON"
  Kokkos_ENABLE_DEPRECATED_CODE_3:BOOL="OFF"
  Kokkos_ENABLE_DEPRECATED_CODE_4:BOOL="OFF"
  Kokkos_ENABLE_OPENMP:BOOL="OFF"
  Kokkos_ENABLE_SERIAL:BOOL="ON"
  Kokkos_ENABLE_TESTS="ON"

This (rebased) PR builds on my machine with -j 16 in 20 mins. Current develop takes 70 mins.

I have tested this on a Pascal arch machine, all tests pass (except for the 1gb memory test since I don't have that on my machine. I specifically tested the random number tests 100 times in a row with no failure.

[crtrott]

Doesn't pass Windows CUDA testing (doesn't like noreturn on something)

[dalg24]

Do we still need

kokkos/core/src/Kokkos_Abort.hpp

Lines 39 to 46 in 835dbf5

	#if defined(KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK)
	// required to workaround failures in random number generator unit tests with
	// pre-volta architectures
	#define KOKKOS_IMPL_ABORT_NORETURN
	#else
	// cuda_abort aborts when building for other platforms than macOS
	#define KOKKOS_IMPL_ABORT_NORETURN [[noreturn]]
	#endif

then?

Can/should we handle the [[noreturn]] part elsewhere?
Isn't this fix mostly about using the __noinline__ specifier?

[dalg24]

OK I can see how the extern specifier didn't make much sense but, out of curiosity, is that change necessary (not asking to change anything at this time).
Also, for consistency I would have slightly preferred if the [[noreturn]] attribute came first.

[nmm0]

not sure if the extern change was necessary -- but looks like clang cuda also chokes on the [[noreturn]] not being first so I moved it

[dalg24]

Retest this please

[crtrott]

Clang CUDA doesn't work:

/var/jenkins/workspace/Kokkos_PR-7013/core/src/Cuda/Kokkos_Cuda_abort.hpp:31:12: error: an attribute list cannot appear here [clang-diagnostic-error]
__device__ [[noreturn]] void __assertfail(const void *message, const void *file,

[nmm0]

Clang CUDA doesn't work:

/var/jenkins/workspace/Kokkos_PR-7013/core/src/Cuda/Kokkos_Cuda_abort.hpp:31:12: error: an attribute list cannot appear here [clang-diagnostic-error]
__device__ [[noreturn]] void __assertfail(const void *message, const void *file,

@crtrott let me see if reversing the order of [[noreturn]] will do it. It's more consistent anyway

[nmm0]

Do we still need

kokkos/core/src/Kokkos_Abort.hpp

Lines 39 to 46 in 835dbf5

	#if defined(KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK)
	// required to workaround failures in random number generator unit tests with
	// pre-volta architectures
	#define KOKKOS_IMPL_ABORT_NORETURN
	#else
	// cuda_abort aborts when building for other platforms than macOS
	#define KOKKOS_IMPL_ABORT_NORETURN [[noreturn]]
	#endif

then?

Can/should we handle the [[noreturn]] part elsewhere? Isn't this fix mostly about using the __noinline__ specifier?

@dalg24 interestingly I get Kokkos_Abort.hpp(99): warning #1305-D: function declared with "noreturn" does return when trying to remove so I guess in some code paths it can return

[nmm0]

Do we still need

kokkos/core/src/Kokkos_Abort.hpp

Lines 39 to 46 in 835dbf5

	#if defined(KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK)
	// required to workaround failures in random number generator unit tests with
	// pre-volta architectures
	#define KOKKOS_IMPL_ABORT_NORETURN
	#else
	// cuda_abort aborts when building for other platforms than macOS
	#define KOKKOS_IMPL_ABORT_NORETURN [[noreturn]]
	#endif

then?
Can/should we handle the [[noreturn]] part elsewhere? Isn't this fix mostly about using the __noinline__ specifier?

@dalg24 interestingly I get Kokkos_Abort.hpp(99): warning #1305-D: function declared with "noreturn" does return when trying to remove so I guess in some code paths it can return

Also, just tested and (maybe unsurprisingly) changing this doesn't change compilation time.

[crtrott]

There are errors with illegal instructions. Is that real?

[nmm0]

There are errors with illegal instructions. Is that real?

I suspect maybe they are... I'm looking at it this week

[nmm0]

Retest this please