[amdgpu] Drop lowering of LDS used by global variables

[JonChesterfield]

Summary

Approximately revert D103431.

LDS variables are allocated at kernel launch and deallocated at kernel exit.
The address is therefore kernel execution dependent. Global variables are
initialized by values written to .data, which can't be done for a LDS variable
as there is no kernel running, or by a global constructor. Initializing the
global to the address of some LDS allocated by a global constructor is possible
but indistinguishable from undef.

Assigning the address of a LDS variable to a global should be a sema error. It
isn't for openmp, haven't checked other languages. Failing that it could be set
to undef, perhaps in this pass.

[JonChesterfield]

Passed internal CI 'psdb', ID 4593.

Deleting this obfuscation unblocks a patch to avoid allocating LDS from kernels that (sufficiently clearly) don't need it.

[ronlieb]

LGTM, but needs Matt or Stas or Brian ...

[JonChesterfield]

• rebase

[JonChesterfield]

Ping. Would like to ship this by the weekend.

[rampitec]

I can see how it is illegal to assign an address of an LDS object to a global pointer global variable. What about a global variable which is an LDS pointer? I.e.:

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

This should be legal I guess.

[JonChesterfield]

Yes, agreed. The address of an LDS variable is scoped to the kernel execution, thus some other LDS variable could be initialised with it. Fortunately, LDS have undef as an initialiser (which is also checked before this call), so that edge case can be postponed until implementing initializers for LDS.

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

123

LDS with non-undef init skipped over here

[JonChesterfield]

• rebase

[arsenm]

I'd rather keep the tests around. The backend needs some stressing of the IR patterns

[JonChesterfield]

The tests of nonsensical IR input being mapped to some other nonsensical IR output? This pass (once fixed) will leave them unchanged, so they'll show nothing

[arsenm]

In D115413#3189737, @JonChesterfield wrote:

The tests of nonsensical IR input being mapped to some other nonsensical IR output? This pass (once fixed) will leave them unchanged, so they'll show nothing

By definition it's not nonsensical IR input since it passes the verifier, and the backend has to have some behavior on it

[JonChesterfield]

I can change the IR tests to ones that run to the backend (instead of just this pass) and see how that behaves. Might crash, might write nonsense to .data. That's strictly an increase in coverage to what we have now, with or without this patch.

It seems unfortunate that we don't have an amdgpu specific verifier pass that rejects invalid addrspace annotated variables.

[rampitec]

In D115413#3189791, @JonChesterfield wrote:

I can change the IR tests to ones that run to the backend (instead of just this pass) and see how that behaves. Might crash, might write nonsense to .data. That's strictly an increase in coverage to what we have now, with or without this patch.

It seems unfortunate that we don't have an amdgpu specific verifier pass that rejects invalid addrspace annotated variables.

Perhaps these tests would be better if globals would have address space 3. When we support LDS initializers this is what we will want to transform.
That said, if we eventually allow LDS initializers isn't that a kind of code we would like to have? What is a motivation to remove it modulo the initializers?

[JonChesterfield]

I am doubtful that this would be the correct strategy for initializers of LDS - I suspect that's going to have to be handled by injecting initialisation code into the prologue of kernels and changing the initialisers to undef - but even if it is, we can retrieve it from git at will. Right now it makes callgraph crawling to avoid unnecessary allocation of LDS much more complicated (or imprecise).

I definitely agree that LDS handling is under-tested at present but there's no value I can see to these specific tests. They'll need to be changed to be more useful at which point we're better off deciding what it is we want to check and writing tests that hit that.

[rampitec]

In D115413#3190617, @JonChesterfield wrote:

I am doubtful that this would be the correct strategy for initializers of LDS - I suspect that's going to have to be handled by injecting initialisation code into the prologue of kernels and changing the initialisers to undef - but even if it is, we can retrieve it from git at will. Right now it makes callgraph crawling to avoid unnecessary allocation of LDS much more complicated (or imprecise).

I definitely agree that LDS handling is under-tested at present but there's no value I can see to these specific tests. They'll need to be changed to be more useful at which point we're better off deciding what it is we want to check and writing tests that hit that.

OK, makes sense.

[JonChesterfield]

Will land this shortly.

I think a significant missing piece of LDS testing is code that runs on the GPU. We have some coarse coverage through the openmp tests (because the openmp runtime happens to make fairly heavy use of LDS) but I'm not aware of any targeted testing. I'm planning to add some through openmp, syntax comes out as

uint64_t var __attribute__((loader_uninitialized));
#pragma allocate(var) allocator(omp_pteam_mem_alloc)

but it's not going to be as precise as it could be through a runtime that doesn't use LDS, will see if I can find some hip or opencl tests that run in tree.

edit: openmp tests are under libomptarget, hip tests are in the llvm-test-suite, e.g. https://reviews.llvm.org/D99997

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);
-
 std::vector<GlobalVariable *> findVariablesToLower(Module &M,
                                                    const Function *F = nullptr);
 
 /// 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 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;
-}
-
-/// \returns true if an LDS global requires 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.
 static 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)) {
-      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);
+    if (isa<GlobalValue>(V)) {
+      // This use of the LDS variable is the initializer of a global variable.
+      // This is ill formed. The address of an LDS variable is kernel dependent
+      // and unknown until runtime. It can't be written to a global variable.
       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;
     }
 
     // User V should be a constant, recursively visit users of V.
     assert(isa<Constant>(V) && "Expected a constant.");
     append_range(Stack, V->users());
   }
 
-  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 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.

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 4
-; 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] }
-
-; 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 2)
-; 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], [1 x float], [2 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 1) 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
-}

llvm/test/CodeGen/AMDGPU/lower-module-lds-indirect.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
-
-; CHECK: %llvm.amdgcn.module.lds.t = type { double, float }
-
-; CHECK: @function_indirect = addrspace(1) global float* addrspacecast (float 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 float*), align 8
-
-; CHECK: @kernel_indirect = addrspace(1) global double* addrspacecast (double addrspace(3)* getelementptr inbounds (%llvm.amdgcn.module.lds.t, %llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds, i32 0, i32 0) to double*), align 8
-
-; CHECK: @llvm.amdgcn.module.lds = internal addrspace(3) global %llvm.amdgcn.module.lds.t undef, align 8
-
-@function_target = addrspace(3) global float undef, align 4
-@function_indirect = addrspace(1) global float* addrspacecast (float addrspace(3)* @function_target to float*), align 8
-
-@kernel_target = addrspace(3) global double undef, align 8
-@kernel_indirect = addrspace(1) global double* addrspacecast (double addrspace(3)* @kernel_target to double*), align 8
-
-; CHECK-LABEL: @function(float %x)
-; CHECK: %0 = load float*, float* addrspace(1)* @function_indirect, align 8
-define void @function(float %x) local_unnamed_addr #5 {
-entry:
-  %0 = load float*, float* addrspace(1)* @function_indirect, align 8
-  store float %x, float* %0, align 4
-  ret void
-}
-
-; CHECK-LABEL: @kernel(double %x)
-; CHECK: call void @llvm.donothing() [ "ExplicitUse"(%llvm.amdgcn.module.lds.t addrspace(3)* @llvm.amdgcn.module.lds) ]
-; CHECK: %0 = load double*, double* addrspace(1)* @kernel_indirect, align 8
-define amdgpu_kernel void @kernel(double %x) local_unnamed_addr #5 {
-entry:
-  %0 = load double*, double* addrspace(1)* @kernel_indirect, align 8
-  store double %x, double* %0, align 8
-  ret void
-}
-
-
-
-