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

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

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@310306 91177308-0d34-0410-b5e6-96231b3b80d8

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)) {
 

test/OpenMP/target_parallel_no_exceptions.cpp

+/// 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
+