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

[jhuber6]

Summary

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.

[jdoerfert]

LG

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

 //===----RTLs/cuda/src/rtl.cpp - Target RTLs Implementation ------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // RTL for CUDA machine
 //
 //===----------------------------------------------------------------------===//
 
 #include <cassert>
 #include <cstddef>
 #include <cuda.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <string>
 #include <unordered_map>
 #include <vector>
 
 #include "Debug.h"
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
         }
 
         if (ExecModeVal < 0 || ExecModeVal > 2) {
           DP("Error wrong exec_mode value specified in cubin file: %d\n",
              ExecModeVal);
           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 "
       }
 
       KernelsList.emplace_back(Func, ExecModeVal);
 
       __tgt_offload_entry Entry = *E;
       Entry.addr = &KernelsList.back();
       addOffloadEntry(DeviceId, Entry);
     }

openmp/libomptarget/test/offloading/global_constructor.cpp

+// 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");
+}