diff options
| -rw-r--r-- | clang/include/clang/Basic/LangOptions.def | 1 | ||||
| -rw-r--r-- | clang/lib/Frontend/CompilerInvocation.cpp | 1 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaExprCXX.cpp | 6 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaStmt.cpp | 6 | ||||
| -rw-r--r-- | clang/test/OpenMP/nvptx_target_exceptions_messages.cpp | 71 |
5 files changed, 83 insertions, 2 deletions
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index f868110719e..06608dda76d 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -202,6 +202,7 @@ LANGOPT(OpenMPSimd , 1, 0, "Use SIMD only OpenMP support.") LANGOPT(OpenMPUseTLS , 1, 0, "Use TLS for threadprivates or runtime calls") LANGOPT(OpenMPIsDevice , 1, 0, "Generate code only for OpenMP target device") LANGOPT(OpenMPCUDAMode , 1, 0, "Generate code for OpenMP pragmas in SIMT/SPMD mode") +LANGOPT(OpenMPHostCXXExceptions , 1, 0, "C++ exceptions handling in the host code.") LANGOPT(RenderScript , 1, 0, "RenderScript") LANGOPT(CUDAIsDevice , 1, 0, "compiling for CUDA device") diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index a6dc1678d1d..ddaa16477aa 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2586,6 +2586,7 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, // Set the flag to prevent the implementation from emitting device exception // handling code for those requiring so. + Opts.OpenMPHostCXXExceptions = Opts.Exceptions && Opts.CXXExceptions; if (Opts.OpenMPIsDevice && T.isNVPTX()) { Opts.Exceptions = 0; Opts.CXXExceptions = 0; diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index a24ae7a077e..95ea8418e1a 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -695,7 +695,11 @@ ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex, bool IsThrownVarInScope) { // Don't report an error if 'throw' is used in system headers. if (!getLangOpts().CXXExceptions && - !getSourceManager().isInSystemHeader(OpLoc)) + !getSourceManager().isInSystemHeader(OpLoc) && + (!getLangOpts().OpenMPIsDevice || + !getLangOpts().OpenMPHostCXXExceptions || + isInOpenMPTargetExecutionDirective() || + isInOpenMPDeclareTargetContext())) Diag(OpLoc, diag::err_exceptions_disabled) << "throw"; // Exceptions aren't allowed in CUDA device code. diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp index c3a627f1d36..d3224b75f45 100644 --- a/clang/lib/Sema/SemaStmt.cpp +++ b/clang/lib/Sema/SemaStmt.cpp @@ -3942,7 +3942,11 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock, ArrayRef<Stmt *> Handlers) { // Don't report an error if 'try' is used in system headers. if (!getLangOpts().CXXExceptions && - !getSourceManager().isInSystemHeader(TryLoc)) + !getSourceManager().isInSystemHeader(TryLoc) && + (!getLangOpts().OpenMPIsDevice || + !getLangOpts().OpenMPHostCXXExceptions || + isInOpenMPTargetExecutionDirective() || + isInOpenMPDeclareTargetContext())) Diag(TryLoc, diag::err_exceptions_disabled) << "try"; // Exceptions aren't allowed in CUDA device code. diff --git a/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp new file mode 100644 index 00000000000..15c9522540f --- /dev/null +++ b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp @@ -0,0 +1,71 @@ +// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -fexceptions -fcxx-exceptions +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fexceptions -fcxx-exceptions -ferror-limit 100 + +#ifndef HEADER +#define HEADER + +template <typename T> +class TemplateClass { + T a; +public: + TemplateClass() { throw 1;} + T f_method() const { return a; } +}; + +int foo(); + +int baz1(); + +int baz2(); + +int baz4() { return 5; } + +template <typename T> +T FA() { + TemplateClass<T> s; + return s.f_method(); +} + +#pragma omp declare target +struct S { + int a; + S(int a) : a(a) { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}} +}; + +int foo() { return 0; } +int b = 15; +int d; +#pragma omp end declare target +int c; + +int bar() { return 1 + foo() + bar() + baz1() + baz2(); } + +int maini1() { + int a; + static long aa = 32; + try { +#pragma omp target map(tofrom \ + : a, b) + { + S s(a); + static long aaa = 23; + a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); + if (!a) + throw "Error"; // expected-error {{cannot use 'throw' with exceptions disabled}} + } + } catch(...) { + } + return baz4(); +} + +int baz3() { return 2 + baz2(); } +int baz2() { +#pragma omp target + try { // expected-error {{cannot use 'try' with exceptions disabled}} + ++c; + } catch (...) { + } + return 2 + baz3(); +} + +#endif // HEADER |
