Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Disable exceptions when used in CUDA code #79

Open
wants to merge 1 commit into
base: master
Choose a base branch
from

Conversation

mkuron
Copy link

@mkuron mkuron commented Oct 6, 2021

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.

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.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

1 participant