[PATCH] D110556: [openmp] Add addrspacecast to getOrCreateIdent

[Jon Chesterfield via Phabricator via llvm-commits]

JonChesterfield created this revision.
JonChesterfield added reviewers: jdoerfert, ye-luo, tianshilei1992, pdhaliwal, ronlieb, gregrodgers.
Herald added subscribers: guansong, hiraditya, yaxunl.
JonChesterfield requested review of this revision.
Herald added subscribers: llvm-commits, openmp-commits, sstefan1.
Herald added projects: OpenMP, LLVM.

Fixes 51982. Minor refactor to remove `return x = y` construct.

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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D110556

Files:
  llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
  openmp/libomptarget/test/offloading/bug51982.c


Index: openmp/libomptarget/test/offloading/bug51982.c
===================================================================
--- /dev/null
+++ 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;
+}
Index: llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
===================================================================
--- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -276,16 +276,19 @@
     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);
+      GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+      GV->setAlignment(Align(8));
+      Ident = GV;
+    }
   }
-  return Builder.CreatePointerCast(Ident, IdentPtr);
+
+  return Builder.CreatePointerBitCastOrAddrSpaceCast(Ident, IdentPtr);
 }
 
 Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(StringRef LocStr) {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D110556.375269.patch
Type: text/x-patch
Size: 1997 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20210927/98c0ce64/attachment.bin>