[Openmp-commits] [openmp] 3247329 - [openmp] Add addrspacecast to getOrCreateIdent

[Jon Chesterfield via Openmp-commits]

Author: Jon Chesterfield
Date: 2021-09-30T21:36:31+01:00
New Revision: 32473291072504f306f8d1f20db460b15e395f7c

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

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

Added: 
    openmp/libomptarget/test/offloading/bug51982.c

Modified: 
    llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

Removed: 
    


################################################################################
diff  --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 94217c632e436..36f6fa68ce242 100644
--- 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
index 0000000000000..4211190ec2ddf
--- /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;
+}