[openmp] Add addrspacecast to getOrCreateIdent

Fixes 51982. Adds a missing CreatePointerCast and allocates a global in
the correct address space.

Test case derived from https://github.com/ROCm-Developer-Tools/aomp/\
blob/aomp-dev/test/smoke/nest_call_par2/nest_call_par2.c by deleting
parts while checking the assertion failure still occurred.

Reviewed By: jdoerfert

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

diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 94217c632e436fb545a7b2594912924075bd64ac..36f6fa68ce24273d69e2911330e516b579a8e9ef 100644 (file)
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -276,15 +276,20 @@ Value *OpenMPIRBuilder::getOrCreateIdent(Constant *SrcLocStr,
     for (GlobalVariable &GV : M.getGlobalList())
       if (GV.getValueType() == OpenMPIRBuilder::Ident && GV.hasInitializer())
         if (GV.getInitializer() == Initializer)
-          return Ident = &GV;
-
-    auto *GV = new GlobalVariable(M, OpenMPIRBuilder::Ident,
-                                  /* isConstant = */ true,
-                                  GlobalValue::PrivateLinkage, Initializer);
-    GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
-    GV->setAlignment(Align(8));
-    Ident = GV;
+          Ident = &GV;
+
+    if (!Ident) {
+      auto *GV = new GlobalVariable(
+          M, OpenMPIRBuilder::Ident,
+          /* isConstant = */ true, GlobalValue::PrivateLinkage, Initializer, "",
+          nullptr, GlobalValue::NotThreadLocal,
+          M.getDataLayout().getDefaultGlobalsAddressSpace());
+      GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+      GV->setAlignment(Align(8));
+      Ident = GV;
+    }
   }
+
   return Builder.CreatePointerCast(Ident, IdentPtr);
 }
 

diff --git a/openmp/libomptarget/test/offloading/bug51982.c b/openmp/libomptarget/test/offloading/bug51982.c
new file mode 100644 (file)
index 0000000..4211190
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/bug51982.c
@@ -0,0 +1,24 @@
+// RUN: %libomptarget-compile-generic -O1 && %libomptarget-run-generic
+// -O1 to run openmp-opt
+
+int main(void) {
+  long int aa = 0;
+
+  int ng = 12;
+  int nxyz = 5;
+
+  const long exp = ng * nxyz;
+
+#pragma omp target map(tofrom : aa)
+  for (int gid = 0; gid < nxyz; gid++) {
+#pragma omp parallel for
+    for (unsigned int g = 0; g < ng; g++) {
+#pragma omp atomic
+      aa += 1;
+    }
+  }
+  if (aa != exp) {
+    return 1;
+  }
+  return 0;
+}