Disable exceptions when used in CUDA code

NVCC supports neither exceptions nor std::terminate in device code, but silently ignores them.
When using Clang to compile CUDA code, "reference to __host__ function" errors are raised when using exceptions or std::terminate.
This patch disables exceptions and uses the __trap intrinsic to deliver consistently correct behavior for both NVCC and Clang.
MPARK_BUILTIN_UNREACHABLE cannot be used here as it again results in "reference to __host__ function" errors.

Author: mkuron

include/mpark/config.hpp

@@ -63,9 +63,9 @@
 #define MPARK_CPP14_CONSTEXPR
 #endif
 
-#if __has_feature(cxx_exceptions) || defined(__cpp_exceptions) || \
+#if (__has_feature(cxx_exceptions) || defined(__cpp_exceptions) || \
     (defined(_MSC_VER) && defined(_CPPUNWIND)) || \
-    defined(__EXCEPTIONS)
+    defined(__EXCEPTIONS)) && !defined(__CUDA_ARCH__)
 #define MPARK_EXCEPTIONS
 #endif
 

include/mpark/variant.hpp

@@ -244,9 +244,15 @@ namespace mpark {
     virtual const char *what() const noexcept override { return "bad_variant_access"; }
   };
 
-  [[noreturn]] inline void throw_bad_variant_access() {
+  [[noreturn]]
+#ifdef __CUDACC__
+  __host__ __device__
+#endif
+  inline void throw_bad_variant_access() {
 #ifdef MPARK_EXCEPTIONS
     throw bad_variant_access{};
+#elif defined(__CUDA_ARCH__)
+    __trap();
 #else
     std::terminate();
     MPARK_BUILTIN_UNREACHABLE;