[AMDGPU] Bit of code clean-up around collection of LDS globals which requires lowering

[hsmhsm]

Have made some cosmetic changes and added comments.

[hsmhsm]

Have done further cosmetic changes.

[JonChesterfield]

Some comments inline.

If this corrects behaviour as described in the comment, it needs tests. I can't see a behaviour change in the diff.

If it's a non functional refactor, I think it makes the code substantially less legible than it was before. Especially dislike using multiple booleans for control flow.

llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp

59

Changing from User to GlobalVariable seems reasonable in itself

61

Why all these booleans? They make the control flow more complicated, but after staring at it for a while, don't appear to do anything different to the 'continue' we had before.

68

Do you know of prior art for this style of commenting? In general I like comments to describe the 'why', and only the 'what' if the implementation is very complicated. I don't think llvm does much of the /* this code will do foo */ do_foo() style.

112

this comment is worse than the equivalent code immediately above since it refers to 3 instead of the named value

115

Moving the predicate (the long sequence of tests) into a separate function seems OK. Should be s/continue/return false/g

136

Changing from none of (use requires lowering) to if (can skip) seems ok, not sure if makes any difference

llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h

25–26

I don't think these benefit from comments. Not sure what 'required' can mean in the context of getAlign. Perhaps the clearer thing to do is put the one line functions in the header directly

[hsmhsm]

In D100594#2694280, @JonChesterfield wrote:

Some comments inline.

If this corrects behaviour as described in the comment, it needs tests. I can't see a behaviour change in the diff.

If it's a non functional refactor, I think it makes the code substantially less legible than it was before. Especially dislike using multiple booleans for control flow.

I am confused myself now. I am not able to recollect what was forced me to change the code here. When, I was working on the changes for D91516, there was some test it forced my to re-arrange the code as I did here. Anyhow, let's hold this patch for now, and come back to it, when D91516 necessitates it.

[hsmhsm]

Fixed review comments by Jon - Removed all the new code changes, and have made a small cosmetic
changes to original code.

[hsmhsm]

Removed redundant comments.

[hsmhsm]

I will be abondaning this patch for now, since there is no much clarity about this patch at the moment. I will reopen a new patch when the need arises with much clarity and needed lit tests.

llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h

 
 namespace llvm {
 
 namespace AMDGPU {
 
 bool isKernelCC(Function *Func);
 
 Align getAlign(DataLayout const &DL, const GlobalVariable *GV);
 
-bool userRequiresLowering(const SmallPtrSetImpl<GlobalValue *> &UsedList,
-                          User *InitialUser);
+bool skipLDSLowering(GlobalVariable *GV,
+                     const SmallPtrSetImpl<GlobalValue *> &UsedList);
+
+bool isLDSLowereringRequired(GlobalVariable *GV,
+                             const SmallPtrSetImpl<GlobalValue *> &UsedList);
 
 std::vector<GlobalVariable *>
 findVariablesToLower(Module &M, const SmallPtrSetImpl<GlobalValue *> &UsedList);
 
+// Collect "llvm.used" and "llvm.compiler.used" global vars.
 SmallPtrSet<GlobalValue *, 32> getUsedList(Module &M);
 
 } // end namespace AMDGPU
 
 } // end namespace llvm
 
 #endif // LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H

llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp

bool isKernelCC(Function *Func) {
   return AMDGPU::isModuleEntryFunctionCC(Func->getCallingConv());
 }
 
 Align getAlign(DataLayout const &DL, const GlobalVariable *GV) {
   return DL.getValueOrABITypeAlignment(GV->getPointerAlignment(DL),
                                        GV->getValueType());
 }
 
-bool userRequiresLowering(const SmallPtrSetImpl<GlobalValue *> &UsedList,
-                          User *InitialUser) {
-  // Any LDS variable can be lowered by moving into the created struct
-  // Each variable so lowered is allocated in every kernel, so variables
-  // whose users are all known to be safe to lower without the transform
-  // are left unchanged.
+// Check if we can skip the lowering for current LDS global `GV` based on its
+// usage.
+//
+// If `GV` is used within non-kernel functions, then, it certainly requires
+// lowering. Else, if it is used as global initializer, then, lowering requires
+// if it is used as an initilizer of normal globals (*NON* llvm.used and
+// llvm.compiler.used globals). All other uses of `GV` can be ignored.
+bool skipLDSLowering(GlobalVariable *GV,
+                     const SmallPtrSetImpl<GlobalValue *> &UsedList) {
   SmallPtrSet<User *, 8> Visited;
-  SmallVector<User *, 16> Stack;
-  Stack.push_back(InitialUser);
+  SmallVector<User *, 16> Stack(GV->users());
 
   while (!Stack.empty()) {
     User *V = Stack.pop_back_val();
     Visited.insert(V);
 
     if (auto *G = dyn_cast<GlobalValue>(V->stripPointerCasts())) {
       if (UsedList.contains(G)) {
         continue;
       }
     }
 
     if (auto *I = dyn_cast<Instruction>(V)) {
       if (isKernelCC(I->getFunction())) {
         continue;
       }
     }
 
     if (auto *E = dyn_cast<ConstantExpr>(V)) {
       for (Value::user_iterator EU = E->user_begin(); EU != E->user_end();
            ++EU) {
         if (Visited.insert(*EU).second) {
           Stack.push_back(*EU);
         }
       }
       continue;
     }
 
-    // Unknown user, conservatively lower the variable
-    return true;
-  }
+    // We *cannot* skip the lowering of LDS global `GV`.
+    return false;
+  }
+
 
+bool isLDSLowereringRequired(GlobalVariable *GV,
+                             const SmallPtrSetImpl<GlobalValue *> &UsedList) {
+  if (GV->getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {
     return false;
   }
 
-std::vector<GlobalVariable *>
-findVariablesToLower(Module &M,
-                     const SmallPtrSetImpl<GlobalValue *> &UsedList) {
-  std::vector<llvm::GlobalVariable *> LocalVars;
-  for (auto &GV : M.globals()) {
-    if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {
-      continue;
-    }
-    if (!GV.hasInitializer()) {
-      // addrspace(3) without initializer implies cuda/hip extern __shared__
-      // the semantics for such a variable appears to be that all extern
-      // __shared__ variables alias one another, in which case this transform
-      // is not required
-      continue;
-    }
-    if (!isa<UndefValue>(GV.getInitializer())) {
-      // Initializers are unimplemented for local address space.
-      // Leave such variables in place for consistent error reporting.
-      continue;
-    }
-    if (GV.isConstant()) {
-      // A constant undef variable can't be written to, and any load is
-      // undef, so it should be eliminated by the optimizer. It could be
-      // dropped by the back end if not. This pass skips over it.
-      continue;
-    }
-    if (std::none_of(GV.user_begin(), GV.user_end(), [&](User *U) {
-          return userRequiresLowering(UsedList, U);
-        })) {
-      continue;
-    }
+  if (!GV->hasInitializer()) {
+    // addrspace(3) without initializer implies cuda/hip extern __shared__ the
+    // semantics for such a variable appears to be that all extern // __shared__
+    // variables alias one another, in which case this transform is not required
+    return false;
+
+  if (!isa<UndefValue>(GV->getInitializer())) {
+    // Initializers are unimplemented for local address space. Leave such
+    // variables in place for consistent error reporting.
+    return false;
+
+  if (GV->isConstant()) {
+    // A constant undef variable can't be written to, and any load is undef, so
+    // it should be eliminated by the optimizer. It could be dropped by the back
+    // end if not. This pass skips over it.
+    return false;
+
+  if (skipLDSLowering(GV, UsedList)) {
+    // We can safely ignore users of GV, hence lowering of `GV` is not required.
+    return false;
+
+  return true;
+
+// Collect all those LDS globals which requires lowering.
+std::vector<GlobalVariable *>
+findVariablesToLower(Module &M,
+                     const SmallPtrSetImpl<GlobalValue *> &UsedList) {
+  std::vector<llvm::GlobalVariable *> LocalVars;
+  for (auto &GV : M.globals()) {
+    if (isLDSLowereringRequired(&GV, UsedList))
       LocalVars.push_back(&GV);
   }
   return LocalVars;
 }
 
+// Collect "llvm.used" and "llvm.compiler.used" global vars.
 SmallPtrSet<GlobalValue *, 32> getUsedList(Module &M) {
   SmallPtrSet<GlobalValue *, 32> UsedList;
 
   SmallVector<GlobalValue *, 32> TmpVec;
   collectUsedGlobalVariables(M, TmpVec, true);
   UsedList.insert(TmpVec.begin(), TmpVec.end());
 
   TmpVec.clear();
   collectUsedGlobalVariables(M, TmpVec, false);
   UsedList.insert(TmpVec.begin(), TmpVec.end());
 
   return UsedList;
 }
 
 } // end namespace AMDGPU
 
 } // end namespace llvm