[OpenMP] Prevent emission of exception handling code when using OpenMP to offload...

Summary: For the OpenMP toolchain which offloads to NVIDIA GPUs make sure that no exception handling code is emitted.

Reviewers: arpith-jacob, sfantao, caomhin, carlo.bertolli, ABataev, Hahnfeld, hfinkel, tstellar

Reviewed By: ABataev, Hahnfeld

Subscribers: rengolin, Hahnfeld, cfe-commits

Differential Revision: https://reviews.llvm.org/D29904

llvm-svn: 310306

diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 003ea55..eb73e28 100644 (file)
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -2308,6 +2308,13 @@ 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.
+  if (Opts.OpenMPIsDevice && T.isNVPTX()) {
+    Opts.Exceptions = 0;
+    Opts.CXXExceptions = 0;
+  }
+
   // Get the OpenMP target triples if any.
   if (Arg *A = Args.getLastArg(options::OPT_fopenmp_targets_EQ)) {
 

diff --git a/clang/test/OpenMP/target_parallel_no_exceptions.cpp b/clang/test/OpenMP/target_parallel_no_exceptions.cpp
new file mode 100644 (file)
index 0000000..95189a3
--- /dev/null
+++ b/clang/test/OpenMP/target_parallel_no_exceptions.cpp
@@ -0,0 +1,18 @@
+/// Make sure no exception messages are inclided in the llvm output.
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -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 - | FileCheck %s --check-prefix CHK-EXCEPTION
+
+void test_increment() {
+#pragma omp target
+  {
+    []() { return; }();
+  }
+}
+
+int main() {
+  test_increment();
+  return 0;
+}
+
+//CHK-EXCEPTION-NOT: invoke
+