[Libomptarget] Correctly default to Generic if exec_mode is not present

Currently, the runtime returns an error when the `exec_mode` global is
not present. The expected behvaiour is that the region will default to
Generic. This prevents global constructors from being called because
they do not contain execution mode globals.

Reviewed By: jdoerfert

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

diff --git a/openmp/libomptarget/plugins/cuda/src/rtl.cpp b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
index f69dc76..44fc672 100644 (file)
--- a/openmp/libomptarget/plugins/cuda/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
@@ -829,10 +829,9 @@ public:
           return nullptr;
         }
       } else {
-        REPORT("Loading global exec_mode '%s' - symbol missing, using default "
-               "value GENERIC (1)\n",
-               ExecModeName);
-        CUDA_ERR_STRING(Err);
+        DP("Loading global exec_mode '%s' - symbol missing, using default "
+           "value GENERIC (1)\n",
+           ExecModeName);
       }
 
       KernelsList.emplace_back(Func, ExecModeVal);

diff --git a/openmp/libomptarget/test/offloading/global_constructor.cpp b/openmp/libomptarget/test/offloading/global_constructor.cpp
new file mode 100644 (file)
index 0000000..be78759
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/global_constructor.cpp
@@ -0,0 +1,19 @@
+// RUN: %libomptarget-compilexx-generic && %libomptarget-run-generic | %fcheck-generic
+
+#include <cmath>
+#include <cstdio>
+
+const double Host = log(2.0) / log(2.0);
+#pragma omp declare target
+const double Device = log(2.0) / log(2.0);
+#pragma omp end declare target
+
+int main() {
+  double X;
+#pragma omp target map(from : X)
+  { X = Device; }
+
+  // CHECK: PASS
+  if (X == Host)
+    printf("PASS\n");
+}