Disable exceptions when used in CUDA code

[mkuron]

CUDA supports neither exceptions nor std::terminate in device code, but NVCC silently ignores them. Unfortunately, Clang, which can be used as a drop-in replacement for NVCC to compile CUDA code, does not silently ignore these things, raising the following errors:

variant/include/mpark/variant.hpp:1811:42: error: reference to __host__ function 'throw_bad_variant_access' in __host__ __device__ function
          holds_alternative<I>(v) ? 0 : (throw_bad_variant_access(), 0))(

variant/include/mpark/variant.hpp:242:9: error: reference to __host__ function '~exception' in __host__ __device__ function
  class bad_variant_access : public std::exception {

To observe this issue, use the following minimal piece of sample code:

#include <mpark/variant.hpp>

using V = mpark::variant<int>;

__global__
void kernel(V v) {
    mpark::get<0>(v);
}

int main() {
    kernel<<<1,1>>>(0);
    return 0;
}

It can be successfully compiled by NVCC (/usr/local/cuda-11.0/bin/nvcc -I variant/include -x cu -ccbin $(which g++-9) -std=c++14 -gencode=arch=compute_52,code=[sm_52,compute_52] --expt-relaxed-constexpr -o test_nvcc test.cpp). It cannot be successfully compiled by Clang (clang++-10 -I variant/include -x cuda --gcc-toolchain=$(dirname $(dirname $(which gcc-9))) -std=c++14 --cuda-path=/usr/local/cuda-10.0 --cuda-gpu-arch=sm_52 -L /usr/local/cuda-10.1/lib -lcudart -o test_clang test.cpp) without my patch. Note that the precise versions of CUDA, Clang or GCC used in the compile command do not matter, the ones I gave here simply are whatever compatible versions I had available on my system.

My patch disables exceptions inside CUDA device code and uses the __trap intrinsic as an std::terminate-equivalent to deliver consistently correct behavior for both NVCC and Clang.