diff --git a/include/mpark/config.hpp b/include/mpark/config.hpp
index 6e2445d89..913ed9713 100644
--- a/include/mpark/config.hpp
+++ b/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
 
diff --git a/include/mpark/variant.hpp b/include/mpark/variant.hpp
index 2fb2ac549..68991365e 100644
--- a/include/mpark/variant.hpp
+++ b/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;