Remove KOKKOS_IMPL_CUDA_* macros

[masterleinad]

Remove most KOKKOS_IMPL_CUDA_* macros. These are internal (and short) and only really have one definition anyway.

[masterleinad]

Retest this please.

[crtrott]

There is one mistake. Otherwise looks ok. But I want another careful review. Screwing this up, like that one mistake, would mostly manifest as super rare race conditions....

[masterleinad]

Retest this please.

[dalg24]

Please elaborate what changes are proposed and why you suggested them. (I don't have an issue with them, just with the way you present them)

[masterleinad]

Please elaborate what changes are proposed and why you suggested them. (I don't have an issue with them, just with the way you present them)

Since requiring CUDA version >9.0, these macros only have one definition and there is no expectation that we need to provide alternatives for other versions. Thus, this pull request is solely about code cleanup removing branches that will never be executed, and avoiding macros that hide the actual code being used. The specific macros being considered are:

• KOKKOS_IMPL_CUDA_ACTIVEMASK
• KOKKOS_IMPL_CUDA_SYNCWARP
• KOKKOS_IMPL_CUDA_SYNCWARP_MASK
• KOKKOS_IMPL_CUDA_BALLOT
• KOKKOS_IMPL_CUDA_BALLOT_MASK
• KOKKOS_IMPL_CUDA_SHFL
• KOKKOS_IMPL_CUDA_SHFL_MASK
• KOKKOS_IMPL_CUDA_SHFL_UP
• KOKKOS_IMPL_CUDA_SHFL_DOWN
• KOKKOS_IMPL_CUDA_SHFL_DOWN_MASK

The other macros in core/src/Cuda/Kokkos_Cuda_Version_9_8_Compatibility.hpp seem useful. The name of the file might be a misnomer now, though.

Additionally, some tests in core/unit_test/TestViewMapping_subview.hpp are reenabled for CUDA >10.0. The intent before was to disable them for CUDA 9 without taking newer versions into account.

[dalg24]

The other macros in core/src/Cuda/Kokkos_Cuda_Version_9_8_Compatibility.hpp seem useful.

Actually we can probably get rid of everything.

KOKKOS_IMPL_CUDA_MAX_SHFL_SIZEOF is not actually used anywhere.

kokkos/core/src/Cuda/Kokkos_Cuda_Vectorization.hpp

Lines 138 to 141 in 1fb0c28

	// TODO DSH shouldn't this be KOKKOS_IMPL_CUDA_MAX_SHFL_SIZEOF instead of
	// sizeof(int)? (Need benchmarks to decide which is faster)
	using shuffle_as_t = int;
	enum : int { N = sizeof(Scalar) / sizeof(shuffle_as_t) };

KOKKOS_IMPL_CUDA_SYNCWARP_OR_RETURN is exclusively used in core/src/Cuda/Kokkos_Cuda_Task.hpp and all occurrences are actually commented except

kokkos/core/src/Cuda/Kokkos_Cuda_Task.hpp

Lines 393 to 395 in 1fb0c28

	#if defined(KOKKOS_ENABLE_DEBUG)
	KOKKOS_IMPL_CUDA_SYNCWARP_OR_RETURN("TaskQueue CUDA task_ptr");
	#endif

.
At the very least we could move the macro definition closer to its point of use.

[ndellingwood]

@masterleinad @dalg24 The KOKKOS_IMPL_CUDA_BALLOT macro was used in Trilinos by Sacado, merge of this PR caused nightly Trilinos builds to fail. Is the correct work around to see about adding macro to Sacado?

#if defined(__CUDA_ARCH__)
#define KOKKOS_IMPL_CUDA_BALLOT(x) __ballot_sync(__activemask(), x)
#else
define KOKKOS_IMPL_CUDA_BALLOT(x) 0
#endif

Edit: This was a build with cuda/9.2 on Kepler architecture

[masterleinad]

Yes, KOKKOS_IMPL_CUDA_BALLOT(x) should be replaced by __ballot_sync(__activemask(), x).