[AMDGPU] Fix missing lowering of LDS used in global scope.

[hsmhsm]

Fixed clang-tidy warning.

[hsmhsm]

Added new lit test-case - lower-kernel-lds-global-uses.ll

[hsmhsm]

Updated the logic within shouldLowerLDSToStruct().

[JonChesterfield]

Is this the proposed fix for the missing lowering of the LDS in the following?

@lds = addrspace(3) global float undef, align 8
@gptr = addrspace(1) global i64* addrspacecast (float addrspace(3)* @lds to i64*), align 8

@llvm.used = appending global [2 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (float addrspace(3)* @lds to i8 addrspace(3)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr to i8 addrspace(1)*) to i8*)], section "llvm.metadata"

define void @f0() {
  %ld = load i64*, i64* addrspace(1)* @gptr
  ret void
}

[hsmhsm]

In D103431#2790962, @JonChesterfield wrote:

Is this the proposed fix for the missing lowering of the LDS in the following?

@lds = addrspace(3) global float undef, align 8
@gptr = addrspace(1) global i64* addrspacecast (float addrspace(3)* @lds to i64*), align 8

@llvm.used = appending global [2 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (float addrspace(3)* @lds to i8 addrspace(3)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr to i8 addrspace(1)*) to i8*)], section "llvm.metadata"

define void @f0() {
  %ld = load i64*, i64* addrspace(1)* @gptr
  ret void
}

yes

[JonChesterfield]

For that test, I note that commenting out the llvm.used statement results in the variable being lowered. I.e. the following works fine:

@lds = addrspace(3) global float undef, align 8
@gptr = addrspace(1) global i64* addrspacecast (float addrspace(3)* @lds to i64*), align 8

define void @f0() {
  %ld = load i64*, i64* addrspace(1)* @gptr
  ret void
}

It's therefore likely that the missed handling is related to skipping uses from the .used list. This change seems to be a complete rewrite of the algorithm in shouldLowerLDSToStruct. I'm not sure we have the test coverage to be confident that a different algorithm will work, can we fix the .used oversight directly instead?

[hsmhsm]

In D103431#2791017, @JonChesterfield wrote:

For that test, I note that commenting out the llvm.used statement results in the variable being lowered. I.e. the following works fine:

@lds = addrspace(3) global float undef, align 8
@gptr = addrspace(1) global i64* addrspacecast (float addrspace(3)* @lds to i64*), align 8

define void @f0() {
  %ld = load i64*, i64* addrspace(1)* @gptr
  ret void
}

It's therefore likely that the missed handling is related to skipping uses from the .used list. This change seems to be a complete rewrite of the algorithm in shouldLowerLDSToStruct. I'm not sure we have the test coverage to be confident that a different algorithm will work, can we fix the .used oversight directly instead?

If the current logic of shouldLowerLDSToStruct() is more robust compere to earlier ones, then, why cannot we use it, instead of patching-up earlier ones?

I spent dedicated time to understand the shortcomings of shouldLowerLDSToStruct() and tried to overcome it in this patch. Please take a look at newly added tests.

[hsmhsm]

Rebased.

[rampitec]

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

64

.equals(). I don't think == does what you want.

122–123

It seems to miss the case when a same constant used by multiple kernels if I am not missing something. We could fix it by converting constant to an instruction but not doing it yet. You even had a utility function for that in some other patch.

130

.equals().

[hsmhsm]

Fixed review comments by Stas.

[hsmhsm]

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

122–123

That is correct - The intention of this patch is to fix bug within shouldLowerLDSToStruct() by making sure that it will properly analyse the uses of LDS to decide if it is to be lowered or not (for both module lds lowering and kernel lds lowering), and also to make the code bit programmer freindly for understanding it.

I implemented a separate patch to handle constants users of LDS for kernel lds lowering (not yet uploaded). I was trying resolve a minor bug where in some cases, LDS was not getting erased from the module. But, now looking at your new patch https://reviews.llvm.org/D103655, looks like I was missing the statement - GV->removeDeadConstantUsers();

But, anyway, let's review the patch https://reviews.llvm.org/D103655.

[rampitec]

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

122–123

Yeah, it took me a good couple of hours to come up with this statement to fix that.

[hsmhsm]

Rebased.

[JonChesterfield]

In D103431#2791057, @hsmhsm wrote:

If the current logic of shouldLowerLDSToStruct() is more robust compere to earlier ones, then, why cannot we use it, instead of patching-up earlier ones?

By current, you mean the code in this patch? It's not obvious to me that this patch is correct. I'm suspicious of branchy code with lots of comments. It looks like other people are already reviewing it, so hopefully they'll catch any other shortcomings.

[foad]

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

64

No, == should be fine. Both == and equals take two StringRefs and do a full string comparison on them.

64

Don't need braces around a single line.

73

A GlobalVariable is a Constant, so the first isa<> check is redundant.

[rampitec]

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

64

I do not see StringRef::operator==() in the documentation?

[hsmhsm]

Fixed review comments by @foad.

[hsmhsm]

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

64

@rampitec Here is the StringRef::operator==() - https://github.com/llvm/llvm-project/blob/main/llvm/include/llvm/ADT/StringRef.h#L897

It infact internally calls .equals() only.

Anyway, I think, it is fine either way except that .equals() is bit verbose, but, I kept .equals() only for now.

[hsmhsm]

Rebased.

[rampitec]

This probably will need rebase on top of D103655 which presumably shall simplify shouldLowerLDSToStruct() logic?

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

64

Thanks!

[hsmhsm]

Rebased, and tried my best to minimize the code changes within shouldLowerLDSToStruct()
based on recent patch https://reviews.llvm.org/D103655.

[hsmhsm]

In D103431#2803681, @rampitec wrote:

This probably will need rebase on top of D103655 which presumably shall simplify shouldLowerLDSToStruct() logic?

Done.

[hsmhsm]

Fixed incorrect and incomplete comment.

[hsmhsm]

UsedList is no longer required to collect "to be lowered" lds.

[hsmhsm]

If the user of LDS is neither global variable nor instruction, then it should be a constant.

[hsmhsm]

explicit "continue" statement at the end of loop is redundant.

[rampitec]

Can you add tests for GlobalAlias? It is not immediately clear if these are handled properly.

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

76

You do not need to cast here, just append U->users().

98–104

stripPointerCasts() still makes sense.

124

No need to cast, you are adding all users anyway. You might want an assert(isa<Constant>(V)) though.

134

How can you get here? You should have inspected all the users by this time and met Instruction uses if any.

[hsmhsm]

Fixed few review comments by Stas.

[hsmhsm]

In D103431#2806206, @rampitec wrote:

Can you add tests for GlobalAlias? It is not immediately clear if these are handled properly.

I have not yet verified the GlobalAlias. Let me look into it.

GlobalAlias is handled and new lit test is added. Please have a look at lower-module-lds-global-alias.ll.

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

98–104

No, as I understand it, we should be very carefull when using stripPointerCasts(), otherwise, it will lead to unexpected and surprising behavior.

For exmaple, in this case, let's try below simple example. Here, only user of "@lds.1" is bitcast inst within kernel. Now, when we are analysing this user, we apply stripPointerCasts(), which will return back "@lds.1" itself, and since it is a global variable, the IF condition satisfies, and we incorrectly land within IF block, where we did not want, which leads to incorrect transormation.

@lds.1 = internal unnamed_addr addrspace(3) global [1 x i8] undef, align 1

define amdgpu_kernel void @k0() {
  %bc = bitcast [1 x i8] addrspace(3)* @lds.1 to i8 addrspace(3)*
  store i8 1, i8 addrspace(3)* %bc, align 1
  ret void
}

134

Please note that the global(s) that we are trying to inspect here are not lds (GV) itself, but, they are the ones which use lds (GV) as an initializer in the global scope or they are special llvm.used. We collect such globals which use lds (GV), and decide whether to lower or not based on the use of those globals.

[hsmhsm]

Added new lit-test for GlobalAlias - lower-module-lds-global-alias.ll.

[rampitec]

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

98–104

Ack.

134

Please note that the global(s) that we are trying to inspect here are not lds (GV) itself, but, they are the ones which use lds (GV) as an initializer in the global scope or they are special llvm.used. We collect such globals which use lds (GV), and decide whether to lower or not based on the use of those globals.

Ack. Makes sense.

[rampitec]

Looks good to me now.

[foad]

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

65

For future reference, SetVector can be useful in cases like this, so that you don't need to declare both a set and a vector and keep them in sync.

134

Could "return true" early here if any G has instruction users, or use std::any_of which will return early for you.

[hsmhsm]

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

134

Correct. Since it does not break any semantics, I will keep it for now, and make suggested code change when I submit any other related patch in near future (we definetely need few more patches sooner than later).

[JonChesterfield]

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

125

After this patch, all users are added. Before it, only those that were not already in the Visited set were added. On the face of it, this seems to discard the guard against cycles. This patch also leaves the Visited set dead (values are inserted into it, but never read).

I don't see any discussion in this patch about removing that guard and it's strange that the set was left in place if that was intentional. Is there a reason cycles cannot occur here?

[JonChesterfield]

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

125

We might be OK - uses can definitely be cyclic, e.g.
uint64_t * __attribute__((address_space(3))) loc = (void*)&loc;
in which case this logic does indeed loop forever, but I don't think it's possible to construct an expression that hits that which gets past the isa<UndefValue> check.

llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp

 
 private:
   bool processUsedLDS(Module &M, Function *F = nullptr) {
     LLVMContext &Ctx = M.getContext();
     const DataLayout &DL = M.getDataLayout();
 
     // Find variables to move into new struct instance
     std::vector<GlobalVariable *> FoundLocalVars =
-        AMDGPU::findVariablesToLower(M, UsedList, F);
+        AMDGPU::findVariablesToLower(M, F);
 
     if (FoundLocalVars.empty()) {
       // No variables to rewrite, no changes made.
       return false;
     }
 
     // Increase the alignment of LDS globals if necessary to maximise the chance
     // that we can use aligned LDS instructions to access them.

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

 class ConstantExpr;
 
 namespace AMDGPU {
 
 bool isKernelCC(const Function *Func);
 
 Align getAlign(DataLayout const &DL, const GlobalVariable *GV);
 
+/// \returns true if a given global variable \p GV (or its global users) appear
+/// as an use within some instruction (either from kernel or from non-kernel).
+bool hasUserInstruction(const GlobalValue *GV);
+
 /// \returns true if an LDS global requres lowering to a module LDS structure
 /// if \p F is not given. If \p F is given it must be a kernel and function
 /// \returns true if an LDS global is directly used from that kernel and it
 /// is safe to replace its uses with a kernel LDS structure member.
-/// \p UsedList contains a union of llvm.used and llvm.compiler.used variables
-/// which do not count as a use.
-bool shouldLowerLDSToStruct(const SmallPtrSetImpl<GlobalValue *> &UsedList,
-                            const GlobalVariable &GV,
+bool shouldLowerLDSToStruct(const GlobalVariable &GV,
                             const Function *F = nullptr);
 
-std::vector<GlobalVariable *>
-findVariablesToLower(Module &M, const SmallPtrSetImpl<GlobalValue *> &UsedList,
+std::vector<GlobalVariable *> findVariablesToLower(Module &M,
                                                    const Function *F = nullptr);
 
 SmallPtrSet<GlobalValue *, 32> getUsedList(Module &M);
 
 /// Replace all uses of constant \p C with instructions in \p F.
 void replaceConstantUsesInFunction(ConstantExpr *C, const Function *F);
 } // end namespace AMDGPU
 
 } // end namespace llvm
 
 #endif // LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H

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

void replaceConstantUsesInFunction(ConstantExpr *C, const Function *F) {
   SetVector<Instruction *> InstUsers;
 
   collectFunctionUses(C, F, InstUsers);
   for (Instruction *I : InstUsers) {
     convertConstantExprsToInstructions(I, C);
   }
 }
 
-bool shouldLowerLDSToStruct(const SmallPtrSetImpl<GlobalValue *> &UsedList,
-                            const GlobalVariable &GV, const Function *F) {
-  // 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.
+bool hasUserInstruction(const GlobalValue *GV) {
+  SmallPtrSet<const User *, 8> Visited;
+  SmallVector<const User *, 16> Stack(GV->users());
+
+  while (!Stack.empty()) {
+    const User *U = Stack.pop_back_val();
+
+    if (!Visited.insert(U).second)
+      continue;
+
+    if (isa<Instruction>(U))
+      return true;
+
+    append_range(Stack, U->users());
+  }
+
+  return false;
+}
+
+bool shouldLowerLDSToStruct(const GlobalVariable &GV, const Function *F) {
+  // We are not interested in kernel LDS lowering for module LDS itself.
+  if (F && GV.getName() == "llvm.amdgcn.module.lds")
+    return false;
+
   bool Ret = false;
   SmallPtrSet<const User *, 8> Visited;
   SmallVector<const User *, 16> Stack(GV.users());
+  SmallPtrSet<const GlobalValue *, 8> GlobalUsers;
 
   assert(!F || isKernelCC(F));
 
   while (!Stack.empty()) {
     const User *V = Stack.pop_back_val();
     Visited.insert(V);
 
-    if (auto *G = dyn_cast<GlobalValue>(V->stripPointerCasts())) {
-      if (UsedList.contains(G)) {
-        continue;
+    if (auto *G = dyn_cast<GlobalValue>(V)) {
+      StringRef GName = G->getName();
+      if (F && GName != "llvm.used" && GName != "llvm.compiler.used") {
+        // For kernel LDS lowering, if G is not a compiler.used list, then we
+        // cannot lower the lds GV since we cannot replace the use of GV within
+        // G.
+        return false;
       }
+      GlobalUsers.insert(G);
+      continue;
     }
 
     if (auto *I = dyn_cast<Instruction>(V)) {
       const Function *UF = I->getFunction();
       if (UF == F) {
         // Used from this kernel, we want to put it into the structure.
         Ret = true;
       } else if (!F) {
+        // For module LDS lowering, lowering is required if the user instruction
+        // is from non-kernel function.
         Ret |= !isKernelCC(UF);
       }
       continue;
     }
 
-    if (auto *E = dyn_cast<ConstantExpr>(V)) {
-      for (const User *U : E->users()) {
-        if (Visited.insert(U).second) {
-          Stack.push_back(U);
-        }
-      }
-      continue;
-    }
+    // User V should be a constant, recursively visit users of V.
+    assert(isa<Constant>(V) && "Expected a constant.");
+    append_range(Stack, V->users());
 
-    // Unknown user, conservatively lower the variable.
-    // For module LDS conservatively means place it into the module LDS struct.
-    // For kernel LDS it means lower as a standalone variable.
-    return !F;
+  if (!F && !Ret) {
+    // For module LDS lowering, we have not yet decided if we should lower GV or
+    // not. Explore all global users of GV, and check if atleast one of these
+    // global users appear as an use within an instruction (possibly nested use
+    // via constant expression), if so, then conservately lower LDS.
+    for (auto *G : GlobalUsers)
+      Ret |= hasUserInstruction(G);
   }
 
   return Ret;
 }
 
-std::vector<GlobalVariable *>
-findVariablesToLower(Module &M, const SmallPtrSetImpl<GlobalValue *> &UsedList,
+std::vector<GlobalVariable *> findVariablesToLower(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 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 (!shouldLowerLDSToStruct(UsedList, GV, F)) {
+    if (!shouldLowerLDSToStruct(GV, F)) {
       continue;
     }
     LocalVars.push_back(&GV);
   }
   return LocalVars;
 }
 
 SmallPtrSet<GlobalValue *, 32> getUsedList(Module &M) {

llvm/test/CodeGen/AMDGPU/lower-kernel-lds-global-uses.ll

+; RUN: opt -S -mtriple=amdgcn-- -amdgpu-lower-module-lds < %s | FileCheck %s
+; RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-module-lds < %s | FileCheck %s
+
+;.
+; Kernel LDS lowering.
+;.
+; @lds.1:  is part of @llvm.used list, and also it is used within kernel, hence it is lowered.
+; @lds.2:  is part of @llvm.compiler.used list, and also it is used within kernel, hence it is lowered.
+; @lds.3:  is used as initializer to @gptr.3, hence @lds.3 is not lowered, though it is used within kernel.
+; @lds.4:  is used as initializer to @gptr.4, hence @lds.4 is not lowered, though it is used within kernel,
+;          irrespective of the uses of @gptr.4 itself ( @gptr.4 is part of llvm.compiler.used list ).
+; @lds.5:  is part of @llvm.used list, but is not used within kernel, hence it is not lowered.
+; @lds.6:  is part of @llvm.compiler.used list, but is not used within kernel, hence it is not lowered.
+;.
+
+; CHECK: %llvm.amdgcn.kernel.k0.lds.t = type { i32, i16 }
+
+; CHECK-NOT: @lds.1
+; CHECK-NOT: @lds.2
+; CHECK: @lds.3 = addrspace(3) global i64 undef, align 8
+; CHECK: @lds.4 = addrspace(3) global float undef, align 4
+; CHECK: @lds.5 = addrspace(3) global i16 undef, align 2
+; CHECK: @lds.6 = addrspace(3) global i32 undef, align 4
+@lds.1 = addrspace(3) global i16 undef, align 2
+@lds.2 = addrspace(3) global i32 undef, align 4
+@lds.3 = addrspace(3) global i64 undef, align 8
+@lds.4 = addrspace(3) global float undef, align 4
+@lds.5 = addrspace(3) global i16 undef, align 2
+@lds.6 = addrspace(3) global i32 undef, align 4
+
+; CHECK: @gptr.3 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* @lds.3 to i64*), align 8
+; CHECK: @gptr.4 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast (float addrspace(3)* @lds.4 to i64 addrspace(3)*) to i64*), align 8
+@gptr.3 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* @lds.3 to i64*), align 8
+@gptr.4 = addrspace(1) global i64* addrspacecast (float addrspace(3)* @lds.4 to i64*), align 8
+
+; CHECK: @llvm.amdgcn.kernel.k0.lds = internal addrspace(3) global %llvm.amdgcn.kernel.k0.lds.t undef, align 4
+
+; CHECK: @llvm.used = appending global [1 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (i16 addrspace(3)* @lds.5 to i8 addrspace(3)*) to i8*)], section "llvm.metadata"
+; CHECK: @llvm.compiler.used = appending global [2 x i8*] [i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.4 to i8 addrspace(1)*) to i8*), i8* addrspacecast (i8 addrspace(3)* bitcast (i32 addrspace(3)* @lds.6 to i8 addrspace(3)*) to i8*)], section "llvm.metadata"
[e-mail] = appending global [2 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (i16 addrspace(3)* @lds.1 to i8 addrspace(3)*) to i8*), i8* addrspacecast (i8 addrspace(3)* bitcast (i16 addrspace(3)* @lds.5 to i8 addrspace(3)*) to i8*)], section "llvm.metadata"
[e-mail] = appending global [3 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (i32 addrspace(3)* @lds.2 to i8 addrspace(3)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.4 to i8 addrspace(1)*) to i8*), i8* addrspacecast (i8 addrspace(3)* bitcast (i32 addrspace(3)* @lds.6 to i8 addrspace(3)*) to i8*)], section "llvm.metadata"
+
+; CHECK-LABEL: @k0()
+; CHECK:   %ld.lds.1 = load i16, i16 addrspace(3)* getelementptr inbounds (%llvm.amdgcn.kernel.k0.lds.t, %llvm.amdgcn.kernel.k0.lds.t addrspace(3)* @llvm.amdgcn.kernel.k0.lds, i32 0, i32 1), align 2
+; CHECK:   %ld.lds.2 = load i32, i32 addrspace(3)* getelementptr inbounds (%llvm.amdgcn.kernel.k0.lds.t, %llvm.amdgcn.kernel.k0.lds.t addrspace(3)* @llvm.amdgcn.kernel.k0.lds, i32 0, i32 0), align 4
+; CHECK:   %ld.lds.3 = load i64, i64 addrspace(3)* @lds.3, align 4
+; CHECK:   %ld.lds.4 = load float, float addrspace(3)* @lds.4, align 4
+; CHECK:   ret void
+define amdgpu_kernel void @k0() {
+  %ld.lds.1 = load i16, i16 addrspace(3)* @lds.1
+  %ld.lds.2 = load i32, i32 addrspace(3)* @lds.2
+  %ld.lds.3 = load i64, i64 addrspace(3)* @lds.3
+  %ld.lds.4 = load float, float addrspace(3)* @lds.4
+  ret void
+}

llvm/test/CodeGen/AMDGPU/lower-module-lds-global-alias.ll

+; RUN: opt -S -mtriple=amdgcn-- -amdgpu-lower-module-lds < %s | FileCheck %s
+; RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-module-lds < %s | FileCheck %s
+
+;.
+; @lds.1:  is aliased with @alias.to.lds.1, and @alias.to.lds.1 is used within kernel @k0.
+;          Hence, @lds.1 is lowered.
+; @lds.2:  is aliased with @alias.to.lds.2, and @alias.to.lds.2 is used within non-kernel @f0,
+;          Hence, @lds.2 is lowered.
+; @lds.3:  is used as initializer to global @gptr.3, and @gptr.3 is aliased with @alias.to.gptr.3,
+;          and @alias.to.gptr.3 is used within kernel @k1. Hence, @lds.3 is lowered.
+; @lds.4:  is used as initializer to global @gptr.4, and @gptr.4 is aliased with @alias.to.gptr.4,
+;          and @alias.to.gptr.4 is used within non-kernel @f1. Hence, @lds.4 is lowered.
+; @lds.5:  is aliased with @alias.to.lds.5, but neither @lds.5 nor @alias.to.lds.5 is used anywhere.
+;          Hence, @lds.5 is not lowered.
+; @lds.6:  is used as initializer to global @gptr.6, and @gptr.6 is aliased with @alias.to.gptr.6.
+;          But none of them are used anywhere. Hence, @lds.6 is not lowered.
+;.
+
+; CHECK: %llvm.amdgcn.module.lds.t = type { [4 x i8], [3 x i8], [1 x i8], [2 x i8], [1 x i8] }
+
+; CHECK-NOT: @lds.1
+; CHECK-NOT: @lds.2
+; CHECK-NOT: @lds.3
+; CHECK-NOT: @lds.4
+; CHECK: @lds.5 = internal unnamed_addr addrspace(3) global [5 x i8] undef, align 8
+; CHECK: @lds.6 = internal unnamed_addr addrspace(3) global [6 x i8] undef, align 8
+@lds.1 = internal unnamed_addr addrspace(3) global [1 x i8] undef, align 1
+@lds.2 = internal unnamed_addr addrspace(3) global [2 x i8] undef, align 2
+@lds.3 = internal unnamed_addr addrspace(3) global [3 x i8] undef, align 4
+@lds.4 = internal unnamed_addr addrspace(3) global [4 x i8] undef, align 4
+@lds.5 = internal unnamed_addr addrspace(3) global [5 x i8] undef, align 8
+@lds.6 = internal unnamed_addr addrspace(3) global [6 x i8] undef, align 8
+
+; CHECK: @gptr.3 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast ([3 x i8] addrspace(3)* getelementptr inbounds (%llvm.amdgcn.module.lds.t, %llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds, i32 0, i32 1) to i64 addrspace(3)*) to i64*), align 8
+; CHECK: @gptr.4 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast (%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds to i64 addrspace(3)*) to i64*), align 8
+; CHECK: @gptr.6 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast ([6 x i8] addrspace(3)* @lds.6 to i64 addrspace(3)*) to i64*), align 8
+@gptr.3 = addrspace(1) global i64* addrspacecast ([3 x i8] addrspace(3)* @lds.3 to i64*), align 8
+@gptr.4 = addrspace(1) global i64* addrspacecast ([4 x i8] addrspace(3)* @lds.4 to i64*), align 8
+@gptr.6 = addrspace(1) global i64* addrspacecast ([6 x i8] addrspace(3)* @lds.6 to i64*), align 8
+
+; CHECK: @llvm.amdgcn.module.lds = internal addrspace(3) global %llvm.amdgcn.module.lds.t undef, align 4
+; CHECK: @llvm.compiler.used = appending global [1 x i8*] [i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds (%llvm.amdgcn.module.lds.t, %llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds, i32 0, i32 0, i32 0) to i8*)], section "llvm.metadata"
+
+; CHECK: @alias.to.lds.1 = alias [1 x i8], getelementptr inbounds (%llvm.amdgcn.module.lds.t, %llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds, i32 0, i32 4)
+; CHECK: @alias.to.lds.2 = alias [2 x i8], getelementptr inbounds (%llvm.amdgcn.module.lds.t, %llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds, i32 0, i32 3)
+; CHECK: @alias.to.gptr.3 = alias i64*, i64* addrspace(1)* @gptr.3
+; CHECK: @alias.to.gptr.4 = alias i64*, i64* addrspace(1)* @gptr.4
+; CHECK: @alias.to.lds.5 = alias [5 x i8], [5 x i8] addrspace(3)* @lds.5
+; CHECK: @alias.to.gptr.6 = alias i64*, i64* addrspace(1)* @gptr.6
[e-mail].1 = alias [1 x i8], [1 x i8] addrspace(3)* @lds.1
[e-mail].2 = alias [2 x i8], [2 x i8] addrspace(3)* @lds.2
[e-mail].3 = alias i64*, i64* addrspace(1)* @gptr.3
[e-mail].4 = alias i64*, i64* addrspace(1)* @gptr.4
[e-mail].5 = alias [5 x i8], [5 x i8] addrspace(3)* @lds.5
[e-mail].6 = alias i64*, i64* addrspace(1)* @gptr.6
+
+; CHECK-LABEL: @f1
+; CHECK:   %ld = load i64*, i64* addrspace(1)* @alias.to.gptr.4, align 8
+; CHECK:   ret void
+define void @f1() {
+  %ld = load i64*, i64* addrspace(1)* @alias.to.gptr.4
+  ret void
+}
+
+; CHECK-LABEL: @f0
+; CHECK:   %bc = bitcast [2 x i8] addrspace(3)* @alias.to.lds.2 to i8 addrspace(3)*
+; CHECK:   store i8 1, i8 addrspace(3)* %bc, align 2
+; CHECK:   ret void
+define void @f0() {
+  %bc = bitcast [2 x i8] addrspace(3)* @alias.to.lds.2 to i8 addrspace(3)*
+  store i8 1, i8 addrspace(3)* %bc, align 2
+  ret void
+}
+
+; CHECK-LABEL: @k1
+; CHECK-LABEL:   call void @llvm.donothing() [ "ExplicitUse"(%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds) ]
+; CHECK-LABEL:   %ld = load i64*, i64* addrspace(1)* @alias.to.gptr.3, align 8
+; CHECK-LABEL:   ret void
+define amdgpu_kernel void @k1() {
+  %ld = load i64*, i64* addrspace(1)* @alias.to.gptr.3
+  ret void
+}
+
+; CHECK-LABEL: @k0
+; CHECK-LABEL:   call void @llvm.donothing() [ "ExplicitUse"(%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds) ]
+; CHECK-LABEL:   %bc = bitcast [1 x i8] addrspace(3)* @alias.to.lds.1 to i8 addrspace(3)*
+; CHECK-LABEL:   store i8 1, i8 addrspace(3)* %bc, align 1
+; CHECK-LABEL:   ret void
+define amdgpu_kernel void @k0() {
+  %bc = bitcast [1 x i8] addrspace(3)* @alias.to.lds.1 to i8 addrspace(3)*
+  store i8 1, i8 addrspace(3)* %bc, align 1
+  ret void
+}

llvm/test/CodeGen/AMDGPU/lower-module-lds-global-uses.ll

+; RUN: opt -S -mtriple=amdgcn-- -amdgpu-lower-module-lds < %s | FileCheck %s
+; RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-module-lds < %s | FileCheck %s
+
+;.
+; @lds.1:  is part of @llvm.used list, and is no-where used. Hence it is not lowered.
+; @lds.2:  is part of @llvm.compiler.used list, and is no-where used. Hence it is not lowered.
+; @lds.3:  is used as initializer to @gptr.3, and is no-where used. @gptr.3 itself is also not
+;          used anywhere else, hence @lds.3 is not lowered.
+; @lds.4:  is used as initializer to @gptr.4, and is no-where used. @gptr.4 is part of
+;          @llvm.compiler.used list, but is no-where else used. hence @lds.4 is not lowered.
+;
+; @lds.5:  is used as initializer to @gptr.5, and is no-where used. @gptr.5 is part of
+;          @llvm.compiler.used list, but is also used within kernel @k0. Hence @lds.5 is lowered.
+; @lds.6:  is used as initializer to @gptr.6, and is no-where used. @gptr.6 is part of
+;          @llvm.compiler.used list, but is also used within non-kernel function @f0. Hence @lds.6 is lowered.
+; @lds.7:  is used as initializer to @gptr.7, and is no-where used. @gptr.7 is used as initializer to @gptr.8,
+;          and @gptr.8 is used within non-kernel function @f1. Hence @lds.7 is lowered.
+;.
+
+; CHECK: %llvm.amdgcn.module.lds.t = type { [3 x float], [4 x i8], [2 x float], [1 x float] }
+
+; CHECK: @lds.1 = addrspace(3) global i16 undef, align 2
+; CHECK: @lds.2 = addrspace(3) global i32 undef, align 4
+; CHECK: @lds.3 = addrspace(3) global i64 undef, align 8
+; CHECK: @lds.4 = addrspace(3) global float undef, align 4
+; CHECK-NOT: @lds.5
+; CHECK-NOT: @lds.6
+; CHECK-NOT: @lds.7
+@lds.1 = addrspace(3) global i16 undef, align 2
+@lds.2 = addrspace(3) global i32 undef, align 4
+@lds.3 = addrspace(3) global i64 undef, align 8
+@lds.4 = addrspace(3) global float undef, align 4
+@lds.5 = addrspace(3) global [1 x float] undef, align 4
+@lds.6 = addrspace(3) global [2 x float] undef, align 8
+@lds.7 = addrspace(3) global [3 x float] undef, align 16
+
+; CHECK: @gptr.3 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* @lds.3 to i64*), align 8
+; CHECK: @gptr.4 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast (float addrspace(3)* @lds.4 to i64 addrspace(3)*) to i64*), align 8
+; CHECK: @gptr.5 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast ([1 x float] addrspace(3)* getelementptr inbounds (%llvm.amdgcn.module.lds.t, %llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds, i32 0, i32 3) to i64 addrspace(3)*) to i64*), align 8
+; CHECK: @gptr.6 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast ([2 x float] addrspace(3)* getelementptr inbounds (%llvm.amdgcn.module.lds.t, %llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds, i32 0, i32 2) to i64 addrspace(3)*) to i64*), align 8
+; CHECK: @gptr.7 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* bitcast (%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds to i64 addrspace(3)*) to i64*), align 8
+; CHECK: @gptr.8 = addrspace(1) global i64** addrspacecast (i64* addrspace(1)* @gptr.7 to i64**), align 8
+@gptr.3 = addrspace(1) global i64* addrspacecast (i64 addrspace(3)* @lds.3 to i64*), align 8
+@gptr.4 = addrspace(1) global i64* addrspacecast (float addrspace(3)* @lds.4 to i64*), align 8
+@gptr.5 = addrspace(1) global i64* addrspacecast ([1 x float] addrspace(3)* @lds.5 to i64*), align 8
+@gptr.6 = addrspace(1) global i64* addrspacecast ([2 x float] addrspace(3)* @lds.6 to i64*), align 8
+@gptr.7 = addrspace(1) global i64* addrspacecast ([3 x float] addrspace(3)* @lds.7 to i64*), align 8
+@gptr.8 = addrspace(1) global i64** addrspacecast (i64* addrspace(1)* @gptr.7 to i64**), align 8
+
+; CHECK: @llvm.used = appending global [1 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (i16 addrspace(3)* @lds.1 to i8 addrspace(3)*) to i8*)], section "llvm.metadata"
+; CHECK: @llvm.amdgcn.module.lds = internal addrspace(3) global %llvm.amdgcn.module.lds.t undef, align 16
+; CHECK: @llvm.compiler.used = appending global [5 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (i32 addrspace(3)* @lds.2 to i8 addrspace(3)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.4 to i8 addrspace(1)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.5 to i8 addrspace(1)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.6 to i8 addrspace(1)*) to i8*), i8* addrspacecast (i8 addrspace(3)* bitcast (%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds to i8 addrspace(3)*) to i8*)], section "llvm.metadata"
[e-mail] = appending global [1 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (i16 addrspace(3)* @lds.1 to i8 addrspace(3)*) to i8*)], section "llvm.metadata"
[e-mail] = appending global [4 x i8*] [i8* addrspacecast (i8 addrspace(3)* bitcast (i32 addrspace(3)* @lds.2 to i8 addrspace(3)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.4 to i8 addrspace(1)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.5 to i8 addrspace(1)*) to i8*), i8* addrspacecast (i8 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.6 to i8 addrspace(1)*) to i8*)], section "llvm.metadata"
+
+; CHECK-LABEL: @f1()
+; CHECK:   %ld = load i64**, i64** addrspace(1)* @gptr.8, align 8
+; CHECK:   ret void
+define void @f1() {
+  %ld = load i64**, i64** addrspace(1)* @gptr.8
+  ret void
+}
+
+; CHECK-LABEL: @f0()
+; CHECK:   %ld = load i32, i32* inttoptr (i64 add (i64 ptrtoint (i32* addrspacecast (i32 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.6 to i32 addrspace(1)*) to i32*) to i64), i64 ptrtoint (i32* addrspacecast (i32
+; CHECK: addrspace(1)* bitcast (i64* addrspace(1)* @gptr.6 to i32 addrspace(1)*) to i32*) to i64)) to i32*), align 4
+; CHECK:   ret void
+define void @f0() {
+  %ld = load i32, i32* inttoptr (i64 add (i64 ptrtoint (i32* addrspacecast (i32 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.6 to i32 addrspace(1)*) to i32*) to i64), i64 ptrtoint (i32* addrspacecast (i32 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.6 to i32 addrspace(1)*) to i32*) to i64)) to i32*), align 4
+  ret void
+}
+
+; CHECK-LABEL: @k0()
+; CHECK:   call void @llvm.donothing() [ "ExplicitUse"(%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds) ]
+; CHECK:   %ld = load i32, i32* inttoptr (i64 add (i64 ptrtoint (i32* addrspacecast (i32 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.5 to i32 addrspace(1)*) to i32*) to i64), i64 ptrtoint (i32* addrspacecast (i32
+; CHECK: addrspace(1)* bitcast (i64* addrspace(1)* @gptr.5 to i32 addrspace(1)*) to i32*) to i64)) to i32*), align 4
+; CHECK:   ret void
+define amdgpu_kernel void @k0() {
+  %ld = load i32, i32* inttoptr (i64 add (i64 ptrtoint (i32* addrspacecast (i32 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.5 to i32 addrspace(1)*) to i32*) to i64), i64 ptrtoint (i32* addrspacecast (i32 addrspace(1)* bitcast (i64* addrspace(1)* @gptr.5 to i32 addrspace(1)*) to i32*) to i64)) to i32*), align 4
+  ret void
+}
+
+; CHECK-LABEL: @k1()
+; CHECK:   call void @llvm.donothing() [ "ExplicitUse"(%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds) ]
+; CHECK:   ret void
+define amdgpu_kernel void @k1() {
+  ret void
+}