[openmp] Add addrspacecast to getOrCreateIdent

[JonChesterfield]

Summary

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.

[JonChesterfield]

@pdhaliwal could you run this locally and confirm it lets you revert the patch to amd-stg-open?

[jdoerfert]

What AS are the "ident globals" generated in if we target amdgcn? If it is not 0 we should generate them in the right AS rather than introducing a cast. If it is 0, this patch looks reasonable.

[JonChesterfield]

In D110556#3024709, @jdoerfert wrote:

What AS are the "ident globals" generated in if we target amdgcn? If it is not 0 we should generate them in the right AS rather than introducing a cast. If it is 0, this patch looks reasonable.

They're in AS1 which seems right. Standard GPU memory. The choice is either we qualify all the deviceRTL functions with an addrspace(1) qualifier, or we create the variable in the GPU memory address space and addrspacecast the pointer to unspecified address space. It works out the same after inlining. This approach follows the pattern we use for stack variables, where they're alloca'ed in AS5 and then immediately addrspacecast to unspecified.

We could remove the AS1 qualification by putting the variable in unspecified address space, but that should lead to either infer address spaces putting it back to AS1 or a minor degradation in codegen (using flat instructions instead of the more specific global one).

[jdoerfert]

We should put it in the right AS by construction. That matches the stack case and is the right thing to do. Casting things after the fact is a necessary evil but not a solution.
I don't know how clang decides they go into AS(1) but we should do the same. Some datalayout or target hook should exist to tell us that. Instead of OpenMPIRBuilder::Ident we
would cast that type first into the AS we need and then create the global. The rest should fall in place.

[JonChesterfield]

The variable is in the right address space when it is returned by new GlobalVariable. This patch doesn't change that. I assume this is derived from the datalayout string but haven't checked that's the mechanism.

The only reason this showed up is the variable in AS1 being passed to a function that didn't specify any address space, which hits the assert in Instruction.cpp.

[JonChesterfield]

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

282

Replacing this constructor call with one that specifies the address space does not solve this bug. It probably makes it more widespread as now all Ident variables would be in AS1 as opposed to only some of them.

Likewise, only casting the newly created GlobalVariable to addrspace(0) does not help, as the one that was returned by return Ident = &GV;

The patch here that ensures the result of getOrCreateIdent is in same address space as IdentPtr (i.e. none) looks necessary regardless of whether the global variable is created in AS1 or ASnone.

I tested this by hardcoding 1

[JonChesterfield]

Thanks for the push to look at this more closely. The change from CreatePointerCast to CreatePointerBitCastOrAddrSpaceCast did nothing. CreatePointerCast ends up calling into the latter. Likewise explicitly putting the variable in the default globals address space doesn't help this bug, but seems a good thing to do anyway.

The actual bug here is at the return Ident = &GV line, which was missed when adding the CreatePointerCast in D84345. The first diff here works because it sends that value through CreatePointerBitCastOrAddrSpaceCast, but an equally valid fix would be to change it to return Ident = Builder.CreatePointerCast(&GV, IdentPtr).

[JonChesterfield]

• Set addrspace on the nearby variable, revert to CreatePointerCast

[jdoerfert]

LG.

[JonChesterfield]

Failed the amdgpu buildbot, in what looks like an independent failure mode. Will try to debug and fix that, then look at re-landing this. Don't really want to land the code change without the test.

[JonChesterfield]

Had to fix two other things first. This test case fails because it is compiled as generic spmd mode, and the amdgpu plugin doesn't know what that is so refuses to load it. Patch at D110845

[JonChesterfield]

Enum patch landed and passed the staging CI, this change works for me locally, going to re-commit. Passed staging CI.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

if (!Ident) {
     Constant *Initializer =
         ConstantStruct::get(OpenMPIRBuilder::Ident, IdentData);
 
     // Look for existing encoding of the location + flags, not needed but
     // minimizes the difference to the existing solution while we transition.
     for (GlobalVariable &GV : M.getGlobalList())
       if (GV.getValueType() == OpenMPIRBuilder::Ident && GV.hasInitializer())
         if (GV.getInitializer() == Initializer)
-          return Ident = &GV;
+          Ident = &GV;
 
-    auto *GV = new GlobalVariable(M, OpenMPIRBuilder::Ident,
-                                  /* isConstant = */ true,
-                                  GlobalValue::PrivateLinkage, Initializer);
+    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);
 }
 
 Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(StringRef LocStr) {
   Constant *&SrcLocStr = SrcLocStrMap[LocStr];
   if (!SrcLocStr) {
     Constant *Initializer =
         ConstantDataArray::getString(M.getContext(), LocStr);

openmp/libomptarget/test/offloading/bug51982.c

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