[Openmp-commits] [openmp] edb8acd - [Libomptarget] Correctly default to Generic if exec_mode is not present

[Joseph Huber via Openmp-commits]

Author: Joseph Huber
Date: 2021-08-18T11:24:28-04:00
New Revision: edb8acdc6ea291308ee3b1b0e697adba183ccd5e

URL: https://github.com/llvm/llvm-project/commit/edb8acdc6ea291308ee3b1b0e697adba183ccd5e
DIFF: https://github.com/llvm/llvm-project/commit/edb8acdc6ea291308ee3b1b0e697adba183ccd5e.diff

LOG: [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

Added: 
    openmp/libomptarget/test/offloading/global_constructor.cpp

Modified: 
    openmp/libomptarget/plugins/cuda/src/rtl.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/plugins/cuda/src/rtl.cpp b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
index f69dc76f615cc..44fc67225d198 100644
--- a/openmp/libomptarget/plugins/cuda/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
@@ -829,10 +829,9 @@ class DeviceRTLTy {
           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
index 0000000000000..be78759d85f9a
--- /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");
+}