From 1d6233649c68bcf324b4fdf1b434486285dbf6b2 Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Wed, 2 May 2018 16:52:07 +0000 Subject: [PATCH] [OPENMP] Enable c++ exceptions outside of the target constructs iff they are enabled for the host. If the compilation for the host enables C++ exceptions, but they are not supported by the device, we still need to allow the code with the exception handling constructs outside of the target regions. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@331372 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/LangOptions.def | 1 + lib/Frontend/CompilerInvocation.cpp | 1 + lib/Sema/SemaExprCXX.cpp | 6 +- lib/Sema/SemaStmt.cpp | 6 +- .../nvptx_target_exceptions_messages.cpp | 71 +++++++++++++++++++ 5 files changed, 83 insertions(+), 2 deletions(-) create mode 100644 test/OpenMP/nvptx_target_exceptions_messages.cpp diff --git a/include/clang/Basic/LangOptions.def b/include/clang/Basic/LangOptions.def index f868110719..06608dda76 100644 --- a/include/clang/Basic/LangOptions.def +++ b/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/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp index a6dc1678d1..ddaa16477a 100644 --- a/lib/Frontend/CompilerInvocation.cpp +++ b/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/lib/Sema/SemaExprCXX.cpp b/lib/Sema/SemaExprCXX.cpp index a24ae7a077..95ea8418e1 100644 --- a/lib/Sema/SemaExprCXX.cpp +++ b/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/lib/Sema/SemaStmt.cpp b/lib/Sema/SemaStmt.cpp index c3a627f1d3..d3224b75f4 100644 --- a/lib/Sema/SemaStmt.cpp +++ b/lib/Sema/SemaStmt.cpp @@ -3942,7 +3942,11 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock, ArrayRef 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/test/OpenMP/nvptx_target_exceptions_messages.cpp b/test/OpenMP/nvptx_target_exceptions_messages.cpp new file mode 100644 index 0000000000..15c9522540 --- /dev/null +++ b/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 +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 +T FA() { + TemplateClass 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(); + 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 -- 2.40.0