From 4fac3739a026702fea7bc804e08705f8a3ced8ce Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Wed, 6 Oct 2021 18:33:39 +0200 Subject: [PATCH] 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. --- include/mpark/config.hpp | 4 ++-- include/mpark/variant.hpp | 8 +++++++- 2 files changed, 9 insertions(+), 3 deletions(-) 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;