Commit a28bbd00 authored by Jon Chesterfield's avatar Jon Chesterfield
Browse files

[amdgpu][nfc] Factor predicate out of findLDSVariablesToLower

parent 824a2881

llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp

+2 −2
Original line number Diff line number Diff line
@@ -158,7 +158,7 @@ public: 


    // 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 @@ public: 
        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.

llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp

+1 −1
Original line number Diff line number Diff line
@@ -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) {

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

+28 −21
Original line number Diff line number Diff line
@@ -105,29 +105,36 @@ static bool shouldLowerLDSToStruct(const GlobalVariable &GV, 
  return Ret;

}



std::vector<GlobalVariable *> findVariablesToLower(Module &M,

                                                   const Function *F) {

  std::vector<llvm::GlobalVariable *> LocalVars;

  for (auto &GV : M.globals()) {

bool isLDSVariableToLower(const GlobalVariable &GV) {

  if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {

      continue;

    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

      continue;

    return false;

  }

  if (!isa<UndefValue>(GV.getInitializer())) {

    // Initializers are unimplemented for LDS address space.

    // Leave such variables in place for consistent error reporting.

      continue;

    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 (!isLDSVariableToLower(GV)) {

      continue;

    }

    if (!shouldLowerLDSToStruct(GV, F)) {

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

+3 −2
Original line number Diff line number Diff line
@@ -29,7 +29,8 @@ namespace AMDGPU { 


Align getAlign(DataLayout const &DL, const GlobalVariable *GV);



std::vector<GlobalVariable *> findVariablesToLower(Module &M,

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.