[llvm] a28bbd0 - [amdgpu][nfc] Factor predicate out of findLDSVariablesToLower

[Jon Chesterfield via llvm-commits]

Author: Jon Chesterfield
Date: 2022-08-31T15:44:51+01:00
New Revision: a28bbd00c65a881bb7073e710d315c23dabbd7c0

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

LOG: [amdgpu][nfc] Factor predicate out of findLDSVariablesToLower

Added: 
    

Modified: 
    llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
    llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp
    llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.cpp
    llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.h

Removed: 
    


################################################################################
diff  --git a/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
index 43d7b99b6cf58..969a98568cc32 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
@@ -158,7 +158,7 @@ class AMDGPULowerModuleLDS : public ModulePass {
 
     // Move variables used by functions into amdgcn.module.lds
     std::vector<GlobalVariable *> ModuleScopeVariables =
-        AMDGPU::findVariablesToLower(M, nullptr);
+        AMDGPU::findLDSVariablesToLower(M, nullptr);
     if (!ModuleScopeVariables.empty()) {
       std::string VarName = "llvm.amdgcn.module.lds";
 
@@ -214,7 +214,7 @@ class AMDGPULowerModuleLDS : public ModulePass {
         continue;
 
       std::vector<GlobalVariable *> KernelUsedVariables =
-          AMDGPU::findVariablesToLower(M, &F);
+          AMDGPU::findLDSVariablesToLower(M, &F);
 
       // Replace all constant uses with instructions if they belong to the
       // current kernel. Unnecessary, removing will cause test churn.

diff  --git a/llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp b/llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp
index aa51c5d20bdc8..c356171f5cc42 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp
@@ -141,7 +141,7 @@ class ReplaceLDSUseImpl {
   std::vector<GlobalVariable *> collectLDSRequiringPointerReplace() {
     // Collect LDS which requires module lowering.
     std::vector<GlobalVariable *> LDSGlobals =
-        llvm::AMDGPU::findVariablesToLower(M, nullptr);
+        llvm::AMDGPU::findLDSVariablesToLower(M, nullptr);
 
     // Remove LDS which don't qualify for replacement.
     llvm::erase_if(LDSGlobals, [&](GlobalVariable *GV) {

diff  --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.cpp
index a5b5be59ac77b..3211c2dbabc3a 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.cpp
@@ -105,29 +105,36 @@ static bool shouldLowerLDSToStruct(const GlobalVariable &GV,
   return Ret;
 }
 
-std::vector<GlobalVariable *> findVariablesToLower(Module &M,
-                                                   const Function *F) {
+bool isLDSVariableToLower(const GlobalVariable &GV) {
+  if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {
+    return false;
+  }
+  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 LDS 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;
+  }
+  return true;
+}
+
+std::vector<GlobalVariable *> findLDSVariablesToLower(Module &M,
+                                                      const Function *F) {
   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 LDS 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.
+    if (!isLDSVariableToLower(GV)) {
       continue;
     }
     if (!shouldLowerLDSToStruct(GV, F)) {

diff  --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.h
index a2d59abd3abb5..53fd4cf335ee0 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.h
@@ -29,8 +29,9 @@ namespace AMDGPU {
 
 Align getAlign(DataLayout const &DL, const GlobalVariable *GV);
 
-std::vector<GlobalVariable *> findVariablesToLower(Module &M,
-                                                   const Function *F);
+bool isLDSVariableToLower(const GlobalVariable &GV);
+std::vector<GlobalVariable *> findLDSVariablesToLower(Module &M,
+                                                      const Function *F);
 
 /// Replace all uses of constant \p C with instructions in \p F.
 void replaceConstantUsesInFunction(ConstantExpr *C, const Function *F);