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

[gtbercea]

Summary

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

[Hahnfeld]

LGTM. Please run clang-format before committing!

test/OpenMP/target_parallel_no_exceptions.cpp

7–8

Not needed, please keep the test as small as possible

[ABataev]

lib/Frontend/CompilerInvocation.cpp

2228–2233

I'm not sure this is the right place for this code.

test/OpenMP/target_parallel_no_exceptions.cpp

5

You can't use #include directive in tests, declare printf function yourself, if you really need it.

[gtbercea]

Clean-up test.

[gtbercea]

lib/Frontend/CompilerInvocation.cpp

2228–2233

Alexey, any suggestion where this code should be moved to?

[Hahnfeld]

Out of interest: This is a limition of the implementation, right? Because the standard only says: A throw executed inside a target region must cause execution to resume within the same target region, and the same thread that threw the exception must catch it.

test/OpenMP/target_parallel_no_exceptions.cpp

15

Please make the test be valid OpenMP: inc is not in declare target and it's also a data race.

I think it's enough to only have a target construct in the test.

[gtbercea]

Redo regression test.

[gtbercea]

lib/Frontend/CompilerInvocation.cpp

2228–2233

Alexey, this code is in the function that initializes Opts.Exceptions and Opts.CXXExceptions and comes after that initialization so the options are not being overwritten by it.

[gtbercea]

run clang-format on test.

[Hahnfeld]

LGTM unless Alexey still has objections

[ABataev]

LG

lib/Frontend/CompilerInvocation.cpp

if (!Opts.OpenMPIsDevice) {
       // Add unsupported host targets here:
       case llvm::Triple::nvptx:
       case llvm::Triple::nvptx64:
         Diags.Report(clang::diag::err_drv_omp_host_target_not_supported)
             << TargetOpts.Triple;
         break;
       }
     }
+
+    // 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)) {
 
     for (unsigned i = 0; i < A->getNumValues(); ++i) {
       llvm::Triple TT(A->getValue(i));
 

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