[AMDGPU] Allocate LDS globals in sorted order of their alignment and size.

[hsmhsm]

Summary

At the start of ISel lowering of kernel function K:

[1] Collect all the LDS globals, say, LDSSet, which are used within the kernel K.
[2] Fix the natural alignment of all the LDS globals from LDSSet.
[3] Sort LDS globals from LDSSet first by their alignment and then by their size order.
[4] Allocate memory for LDS globals from LDSSet in the above sorted order.

[foad]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

20

Isn't this llvm::is_contained?

31

Isn't there a suitable Alignment::max?

92–93

Why is sorting by size important?

I can see that sorting by alignment, descending, avoids alignment gaps (at least in the common case where an object's size is a multiple of its alignment). But sorting by size first means that most globals won't be sorted by alignment (unless they happened to have the same size and different alignments), so you won't get that benefit.

[hsmhsm]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

31

I could find any such function. I could see below ones, but one of the arg should be of MaybeAlign.

inline Align max(MaybeAlign Lhs, Align Rhs) {
  return Lhs && *Lhs > Rhs ? *Lhs : Rhs;
}

inline Align max(Align Lhs, MaybeAlign Rhs) {
  return Rhs && *Rhs > Lhs ? *Rhs : Lhs;
}

92–93

I think we need to first sort based on "alignment", then on tie, based on size.

[rampitec]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

92–93

Probably yes. Strategy to sort by size was used along with the super-alignment, so did the same thing.

[hsmhsm]

Fixed review comments by Jay.

[rampitec]

We probably need a test which only runs this lowering and showing a resulting transformed IR. Test the alignment, sorting and situation with multiple kernels using intersecting LDS variable sets.

[rampitec]

llvm/include/llvm/Support/Alignment.h

348

It does not follow the style around, make the body on separate lines.

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

41

Is there a reason to use std::set and not one of llvm's collections?

52

Can a GV hold a pointer to this GV?

61

It can probably be Operator.

70

Why std::vector?

82

Looks like you are traversing all users just to check if a GC used from this kernel. This is wasteful, you can just create a function checking that a GV is reachable from this kernel and early return if it is.

149

Get DL outside of the loop.

165

This comment shall not be relevant anymore?

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

125

This is mostly a copy-paste of findVariablesToLower() above. I would suggest to factor it somehow.

[foad]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

92–93

I still don't understand if there's a reason to sort by size or name here. It's a stable sort, so why not just sort by alignment?

[rampitec]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

92–93

Size may leave samller gaps. Name is an overkill, merely a backdoor to hack it and force non-trivial "optimizations" by renaming variables. But there are more conceptual issues so far.

[hsmhsm]

Fixed review comments by Stas.

[hsmhsm]

In D102401#2757441, @rampitec wrote:

We probably need a test which only runs this lowering and showing a resulting transformed IR. Test the alignment, sorting and situation with multiple kernels using intersecting LDS variable sets.

Added test - lds-allocation.ll. However, note that we actually do not update the alignment of LDS globals within the module, during allocation, we only just compute the required natural alignment for proper allocation.

llvm/include/llvm/Support/Alignment.h

348

This was forced by clang-format tool.

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

61

I did not understand this comment.

[JonChesterfield]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

92–93

This has been copy/pasted from lowermodulelds, where sort by name was mostly done for more predictable test cases.

[hsmhsm]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

92–93

I thought that it would be helpful in following scenarios. Here we probably know that allocation order is [A, B, C] since both size and alignment are same for all three globals.

@A = internal unnamed_addr addrspace(3) global [4 x i8] undef, align 4
@B = internal unnamed_addr addrspace(3) global [4 x i8] undef, align 4
@C= internal unnamed_addr addrspace(3) global [4 x i8] undef, align 4

[rampitec]

In D102401#2759276, @hsmhsm wrote:

In D102401#2757441, @rampitec wrote:

We probably need a test which only runs this lowering and showing a resulting transformed IR. Test the alignment, sorting and situation with multiple kernels using intersecting LDS variable sets.

Added test - lds-allocation.ll. However, note that we actually do not update the alignment of LDS globals within the module, during allocation, we only just compute the required natural alignment for proper allocation.

I mean an IR test where we can see what lowering has done, not just total allocation amount.

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

61

llvm::Operator.

83

SmallVectorImpl<>

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

79

Maybe just pass a filter lambda there?

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

34

SmallVectorImpl<>.

[arsenm]

llvm/include/llvm/Support/Alignment.h

348

regular std::max works for this

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

21

You're ignoring the explicit alignment of the global, but also implicitly adding the ABI alignment of the type

73–74

cast<>

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

37

I think GV directly has a getAddressSpace

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

34

That doesn't work for a return value

[arsenm]

Can this go into an analysis pass or something? It seems unfortunate that every single function is going to end up doing this same walk through the module

[hsmhsm]

Partially fixed review comments by Matt and Stas.

[hsmhsm]

In D102401#2760924, @arsenm wrote:

Can this go into an analysis pass or something? It seems unfortunate that every single function is going to end up doing this same walk through the module

The only same walk through required is for collecting LDS static globals. If we want to avoid it, then I am not sure, how an analysis result (via an OPT analysis pass) can be used within ISEL stage? Do you know it?

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

21

I think in that case, we need a separate OPT transformation pass just before ISEL (llc) pass, which actually update the natural alignment of LDS globals based on their size, when there is an underaligned case.

61

llvm::Operator is an abstracrtion over Instruction and ConstExpr. An llvm::Operator should be either an Instruction OR ConstExpr, and hence explicit handling of llvm::Operator is not required.

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

79

I personally think that it will be over complicated - collectStaticLDSGlobals() is called by both findVariablesToLower() and getSortedUsedLDSGlobals(). Further filtering of LDS globals is required for findVariablesToLower(), but not for getSortedUsedLDSGlobals().

[hsmhsm]

In D102401#2760317, @rampitec wrote:

In D102401#2759276, @hsmhsm wrote:

In D102401#2757441, @rampitec wrote:

We probably need a test which only runs this lowering and showing a resulting transformed IR. Test the alignment, sorting and situation with multiple kernels using intersecting LDS variable sets.

Added test - lds-allocation.ll. However, note that we actually do not update the alignment of LDS globals within the module, during allocation, we only just compute the required natural alignment for proper allocation.

I mean an IR test where we can see what lowering has done, not just total allocation amount.

Probably then, we need to dump the MIR result by passing the option - "--print-after-isel", and then we need to test for the memory address. But, I see this address is still based on alignment info avaialble in the IR. So, I think, we need to first explicitly update the alignment of LDS globals in a seperate OPT pass just before ISEL pass?

[hsmhsm]

Updated the llvm lit test to check the output of "instruction selection" phasse.

[hsmhsm]

In D102401#2761858, @hsmhsm wrote:

In D102401#2760317, @rampitec wrote:

In D102401#2759276, @hsmhsm wrote:

In D102401#2757441, @rampitec wrote:

We probably need a test which only runs this lowering and showing a resulting transformed IR. Test the alignment, sorting and situation with multiple kernels using intersecting LDS variable sets.

Added test - lds-allocation.ll. However, note that we actually do not update the alignment of LDS globals within the module, during allocation, we only just compute the required natural alignment for proper allocation.

I mean an IR test where we can see what lowering has done, not just total allocation amount.

Probably then, we need to dump the MIR result by passing the option - "--print-after-isel", and then we need to test for the memory address. But, I see this address is still based on alignment info avaialble in the IR. So, I think, we need to first explicitly update the alignment of LDS globals in a seperate OPT pass just before ISEL pass?

I have updated the lit test to check the output of the "instruction selection" pass.

[hsmhsm]

Updated failing lit test - CodeGen/MIR/AMDGPU/machine-function-info.ll

[arsenm]

If we already have a pass that condenses the LDS globals into a single variable to access, what is the advantage of this? Why can't we just always do that compacting and codegen will then not have to worry about optimal LDS packing since it will only see the one global

[hsmhsm]

In D102401#2763394, @arsenm wrote:

If we already have a pass that condenses the LDS globals into a single variable to access, what is the advantage of this? Why can't we just always do that compacting and codegen will then not have to worry about optimal LDS packing since it will only see the one global

But, do we really have one right now? As you might know - "LowerModuleLDSPass" will not guarentee it (atleast as it exist today). We need have something working till we have a kind of perfect solution(s). And, I strongly believe that this patch is really useful patch in the sense it will definetely increase the probability of unaligned ds access to be aligned at runtime. Morever, it only touches kernels, not every single function in the module as you said in one of the earlier patch. I do not see any harm in having this patch, unless I am missing something.

[arsenm]

In D102401#2763512, @hsmhsm wrote:

In D102401#2763394, @arsenm wrote:

If we already have a pass that condenses the LDS globals into a single variable to access, what is the advantage of this? Why can't we just always do that compacting and codegen will then not have to worry about optimal LDS packing since it will only see the one global

But, do we really have one right now? As you might know - "LowerModuleLDSPass" will not guarentee it (atleast as it exist today). We need have something working till we have a kind of perfect solution(s). And, I strongly believe that this patch is really useful patch in the sense it will definetely increase the probability of unaligned ds access to be aligned at runtime. Morever, it only touches kernels, not every single function in the module as you said in one of the earlier patch. I do not see any harm in having this patch, unless I am missing something.

I'm saying you're putting a nontrivial IR inspection into a codegen pass. It would be cleaner if we could leave the allocation optimizations entirely in an IR pass.

[hsmhsm]

In D102401#2763842, @arsenm wrote:

In D102401#2763512, @hsmhsm wrote:

In D102401#2763394, @arsenm wrote:

If we already have a pass that condenses the LDS globals into a single variable to access, what is the advantage of this? Why can't we just always do that compacting and codegen will then not have to worry about optimal LDS packing since it will only see the one global

But, do we really have one right now? As you might know - "LowerModuleLDSPass" will not guarentee it (atleast as it exist today). We need have something working till we have a kind of perfect solution(s). And, I strongly believe that this patch is really useful patch in the sense it will definetely increase the probability of unaligned ds access to be aligned at runtime. Morever, it only touches kernels, not every single function in the module as you said in one of the earlier patch. I do not see any harm in having this patch, unless I am missing something.

I'm saying you're putting a nontrivial IR inspection into a codegen pass. It would be cleaner if we could leave the allocation optimizations entirely in an IR pass.

At the IR level (in the opt pass), you can probably do following:

(1) infer and fix LDS alignment
(2) update the align attribute of LDS globals and LDS memory operations based on above (1)

But, sorting of per kernel used LDS globals can be achieved only in codegen pass as did in this patch. The reason for this is - I do not see any neat way to pass some analysis result which is built in IR pass on top of IR to codegen pass.

[hsmhsm]

In D102401#2763869, @hsmhsm wrote:

In D102401#2763842, @arsenm wrote:

In D102401#2763512, @hsmhsm wrote:

In D102401#2763394, @arsenm wrote:

If we already have a pass that condenses the LDS globals into a single variable to access, what is the advantage of this? Why can't we just always do that compacting and codegen will then not have to worry about optimal LDS packing since it will only see the one global

But, do we really have one right now? As you might know - "LowerModuleLDSPass" will not guarentee it (atleast as it exist today). We need have something working till we have a kind of perfect solution(s). And, I strongly believe that this patch is really useful patch in the sense it will definetely increase the probability of unaligned ds access to be aligned at runtime. Morever, it only touches kernels, not every single function in the module as you said in one of the earlier patch. I do not see any harm in having this patch, unless I am missing something.

I'm saying you're putting a nontrivial IR inspection into a codegen pass. It would be cleaner if we could leave the allocation optimizations entirely in an IR pass.

At the IR level (in the opt pass), you can probably do following:

(1) infer and fix LDS alignment
(2) update the align attribute of LDS globals and LDS memory operations based on above (1)

But, sorting of per kernel used LDS globals can be achieved only in codegen pass as did in this patch. The reason for this is - I do not see any neat way to pass some analysis result which is built in IR pass on top of IR to codegen pass.

May be, we can have it for the time-being, and as a next step, see what best we can do for it.

[hsmhsm]

Moved few utility functions to AMDGPULDSUtils.cpp file since may land-up reusing them in near future.

[arsenm]

In D102401#2763869, @hsmhsm wrote:

(1) infer and fix LDS alignment
(2) update the align attribute of LDS globals and LDS memory operations based on above (1)

But, sorting of per kernel used LDS globals can be achieved only in codegen pass as did in this patch. The reason for this is - I do not see any neat way to pass some analysis result which is built in IR pass on top of IR to codegen pass.

If you rewrite all the accesses relative to a single LDS global, there's nothing for codegen to sort.

[hsmhsm]

In D102401#2763893, @arsenm wrote:

In D102401#2763869, @hsmhsm wrote:

(1) infer and fix LDS alignment
(2) update the align attribute of LDS globals and LDS memory operations based on above (1)

But, sorting of per kernel used LDS globals can be achieved only in codegen pass as did in this patch. The reason for this is - I do not see any neat way to pass some analysis result which is built in IR pass on top of IR to codegen pass.

If you rewrite all the accesses relative to a single LDS global, there's nothing for codegen to sort.

But, do we have any immediate available neat solution(s) which will make sure that all the LDS globals which are used within a kernel is replaced by single struct (or something else) instance? We are struggling with it for a long. Probably, if we solve that problem in near future, then this patch becomes kind of irrelavent and we can revert it at that point, until then, this patch is essential for the immediate need. But, probably let's move all the utility code from AMDGPUMachineFunction to utility file.

[hsmhsm]

Move utility functions from AMDGPUMachineFunction to utility file

[arsenm]

In D102401#2763925, @hsmhsm wrote:

But, do we have any immediate available neat solution(s) which will make sure that all the LDS globals which are used within a kernel is replaced by single struct (or something else) instance?

That's essentially what AMDGPULowerModuleLDSPass does, and could be extended to have a kernel specific table

[rampitec]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

61

No, it is a separate subclass of User: https://llvm.org/doxygen/classllvm_1_1User.html

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

79

I do not like the idea of allocating memory, then copying.

llvm/test/CodeGen/AMDGPU/lds-allocation.ll

35 ↗

(On Diff #345921)

The test checking mir is not any better than a test checking instructions. In fact it is worse because it is using unstable mir interface. It still does not check for an actual layout in any transparent way, it checks the offsets. You can check the offsets in the produces ISA.

It does not help that checks are autogenerated, it obscures relevant checks. For that sort of test you need to check the offsets, nothing more. It also depends on an assumption that stores will not be rescheduled and there is no such guarantee in this case. To w/a it you can store different values (e.g. 1, 2, 3, 4 etc), check that in the ISA, and mark stores volatile.

One thing it does not check is that we do not allocate more LDS than used. It does not check the situation with multiple kernels using intersecting LDS variable sets.

I was looking at a test which will really check memory layout, like that is possible with AMDGPULowerModuleLDSPass. In fact I am probably second to Matt we can try to use AMDGPULowerModuleLDSPass for kernels too, not to multiply entities.

[hsmhsm]

In D102401#2764143, @arsenm wrote:

In D102401#2763925, @hsmhsm wrote:

But, do we have any immediate available neat solution(s) which will make sure that all the LDS globals which are used within a kernel is replaced by single struct (or something else) instance?

That's essentially what AMDGPULowerModuleLDSPass does, and could be extended to have a kernel specific table

Ok, assume that we will extend AMDGPULowerModuleLDSPass to have a kernel specific table. What I am not getting is - how to use this table in order to replace LDS globals by their offset counter-parts within non-kernel functions. Passing kernel-id as an argument is ruled out, since the presence of indirect calls will really break this since function pointers can be used in so many varieties of ways.

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

61

But, if I look into the implemenation of Operator at - https://llvm.org/doxygen/Operator_8h_source.html it clearly tells me that an instance of Operator class will never ever be created, and it is really a place holder for Instrunction and ConstExpr.

/// Return the opcode for this Instruction or ConstantExpr.
unsigned getOpcode() const {
  if (const Instruction *I = dyn_cast<Instruction>(this))
    return I->getOpcode();
  return cast<ConstantExpr>(this)->getOpcode();
}

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

79

Then you mean to say, let's return vector, instead of passing it as reference arg? then we cannot use SmallVectorImpl<> as return type, so, we need to go back to returning SmallVector<> itself.

llvm/test/CodeGen/AMDGPU/lds-allocation.ll

35 ↗

(On Diff #345921)

I think, all of you are kind of suggesting to extend AMDGPULowerModuleLDSPass. But, I am not sure, how it is easily extendible, moreever, it is being invoked before promoteAlloca pass at the moment. We really need to discuss it and come to a common conclusion, before proceeding further.

For me, it looks like - there is a ball, half of which is painted white, and other half black, few are looking at white part, and others looking at black part, and not able to come to a common conclusion.

[foad]

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

79

To avoid allocating memory for a copy of the data, you could filter StaticLDSGlobals in-place with erase(remove_if(StaticLDSGlobals.begin(), StaticLDSGlobals.end(), predicate), StaticLDSGlobals.end()).

[hsmhsm]

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

79

Probably, we land-up with code like below. We still need to copy to LocalVars, because findVariablesToLower() expect return std::vector. And, based on earlier reveiw comments, I have made changes to collectStaticLDSGlobals() to work with SmallVector. Additionally, look at the complexity of the code change in terms of understanding it. So, I thought it is not worth, since we still need to copy anyway.

StaticLDSGlobals.erase(
  std::remove_if(StaticLDSGlobals.begin(), StaticLDSGlobals.end(),
               [&](const GlobalVariable *GV) {
                 GlobalVariable *GV2 = const_cast<GlobalVariable *>(GV);
                  if (std::none_of(GV2->user_begin(), GV2->user_end(), [&](User *U) {
                        return userRequiresLowering(UsedList, U);
                      })) {
                    return true;
                  }
                  return false;
               }),
  StaticLDSGlobals.end());

std::vector<llvm::GlobalVariable *> LocalVars(StaticLDSGlobals.begin(),
                                              StaticLDSGlobals.end());

[hsmhsm]

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

79

Additionally, there is also a added complexity of handling "constness" of "GlobalVariable *"

[foad]

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

79

if (std::none_of(GV2->user_begin(), GV2->user_end(), [&](User *U) {
      return userRequiresLowering(UsedList, U);
    })) {
  return true;
}
return false;

This should just return std::none_of(...);

Additionally, there is also a added complexity of handling "constness" of "GlobalVariable *"

Then fix userRequiresLowering first? Something like: https://reviews.llvm.org/differential/diff/346095/

[hsmhsm]

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

79

This should just return std::none_of(...);

Agree

Then fix userRequiresLowering first? Something like: https://reviews.llvm.org/differential/diff/346095/

I assume, it is not as simple as that, and I guess, it will force me to change within AMDGPULowerModuleLDSPass.cpp which I want to avoid at this point. These kind of changes are better to implement as a kind of code refactor patch. But, nevetheless, if you guys think that it is must, then, I will do it.

[rampitec]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

61

Example:

%sunkaddr4 = getelementptr inbounds i8, i8* bitcast ({ %union, [2000 x i8] }* @global_pointer to i8*), i64 %sunkaddr3

llvm/test/CodeGen/AMDGPU/lds-allocation.ll

35 ↗

(On Diff #345921)

PromoteAlloca is probably not an issue, it will just add a new allocation past that struct. It should be already properly aligned as we are creating it ourselves. The current waste of LDS by the AMDGPULowerModuleLDSPass is an issue though. Ideally it would be nice to come up with a single solution for functions and kernels which I guess we are continue discussing...

[hsmhsm]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

61

If I understand it correctly, @global_pointer is used within a constant expression - bitcast, and this constant expression is used within an instruction - getelementptr.

And this case is neatly handled in the code. I am not sure how relevant it to the "Operator" that you brought up. I still believe, there is no need to check for "Operator".

llvm/test/CodeGen/AMDGPU/lds-allocation.ll

35 ↗

(On Diff #345921)

I still do not understand, if we can really handle it within AMDGPULowerModuleLDSPass. Expecting a single struct variable and only single struct variable within every kernel is too much ideal expectation at this point. And not sure, why you guys are insistiing about it.

Strictly speaking, this pass does not update align of LDS globals? It only fix alignment and allocate memory accordingly. So, the lit tests should check for allocated boundary and total size allocated. And that is what I did here. If you are not happy with MIR result, then, I can check it on final assembly.

To answer some of your earlier comments:

The test checking mir is not any better than a test checking instructions. In fact it is worse because it is using unstable mir interface. It still does not check for an actual layout in any transparent way, it checks the offsets. You can check the offsets in the produces ISA.

I can change it to check instructions instead of mir. Based on the sorting algorithm implemented - we expect which LDS global should be allocated on which assdresses, and how much LDS should be allocated in total. That is what I did here.

It does not help that checks are autogenerated, it obscures relevant checks. For that sort of test you need to check the offsets, nothing more. It also depends on an assumption that stores will not be rescheduled and there is no such guarantee in this case. To w/a it you can store different values (e.g. 1, 2, 3, 4 etc), check that in the ISA, and mark stores volatile.

Let me see, what is missing here.

One thing it does not check is that we do not allocate more LDS than used.

I am checking for it as in "SDAG-HSA: amdhsa_group_segment_fixed_size 31"

It does not check the situation with multiple kernels using intersecting LDS variable sets.

I will try to improve the test cases..

I was looking at a test which will really check memory layout like that is possible with AMDGPULowerModuleLDSPass

Checking which lds global should be allocated on which memory is not memory layout? Could you please give an example of what are you expecting?

In fact I am probably second to Matt we can try to use AMDGPULowerModuleLDSPass for kernels too, not to multiply entities.

I doubt if we can quickly achieve it considering all the mess happening right now.

[hsmhsm]

Rebased to latest main.

[hsmhsm]

Update lit tests as below:

[1] Check instructions instead of MIR
[2] Add test to check LDS used with multiple kernels
[3] Add test to check that llvm.amdgcn.module.lds is allocated on address 0

[hsmhsm]

Fixed comment within lit test.

[hsmhsm]

Within lit tests make sure that different values are stored to LDS globals so that it will be
bit easier to track which memory is allocated for which LDS global.

[rampitec]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

61

OK, I was unable to exploit it, it indeed cast to Constant.

llvm/test/CodeGen/AMDGPU/GlobalISel/lds-allocation.ll

35 ↗

(On Diff #346429)

One remaining problem with these tests now they depend on the register allocation and actual register names. It will be difficult to maintain.

llvm/test/CodeGen/AMDGPU/lds-allocation.ll

35 ↗

(On Diff #345921)

I was looking at a test which will really check memory layout like that is possible with AMDGPULowerModuleLDSPass

Checking which lds global should be allocated on which memory is not memory layout? Could you please give an example of what are you expecting?

I was thinking about a test just showing an IR globals after the transformation, but I guess since it is still just the lowering the IR is not updated and such test is not possible.

In fact I am probably second to Matt we can try to use AMDGPULowerModuleLDSPass for kernels too, not to multiply entities.

I doubt if we can quickly achieve it considering all the mess happening right now.

Given it does not solve minimal per-kernel allocation as we discussed yes, we cannot just use it right away.

[hsmhsm]

Since we are planning different solutions, I am abandoning this patch.

llvm/include/llvm/Support/Alignment.h

 inline Align max(MaybeAlign Lhs, Align Rhs) {
   return Lhs && *Lhs > Rhs ? *Lhs : Rhs;
 }
 
 inline Align max(Align Lhs, MaybeAlign Rhs) {
   return Rhs && *Rhs > Lhs ? *Rhs : Lhs;
 }
 
+inline Align max(Align Lhs, Align Rhs) { return Lhs > Rhs ? Lhs : Rhs; }
+
 #ifndef NDEBUG
 // For usage in LLVM_DEBUG macros.
 inline std::string DebugStr(const Align &A) {
   return std::to_string(A.value());
 }
 // For usage in LLVM_DEBUG macros.
 inline std::string DebugStr(const MaybeAlign &MA) {
   if (MA)

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

 //===-- AMDGPUMachineFunctionInfo.cpp ---------------------------------------=//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "AMDGPUMachineFunction.h"
 #include "AMDGPUPerfHintAnalysis.h"
 #include "AMDGPUSubtarget.h"
+#include "Utils/AMDGPULDSUtils.h"
 #include "llvm/CodeGen/MachineModuleInfo.h"
 #include "llvm/Target/TargetMachine.h"
 
 using namespace llvm;
 
+// Fix natural alignment for LDS global `GV` based on its size.
+static Align fixAlignment(const DataLayout &DL, const GlobalVariable *GV,
+                          Align Alignment) {
+  TypeSize GVSize = DL.getTypeAllocSize(GV->getValueType());
+
+  if (GVSize > 8) {
+    // We might want to use a b96 or b128 load/store
+    Alignment = max(Alignment, Align(16));
+  } else if (GVSize > 4) {
+    // We might want to use a b64 load/store
+    Alignment = max(Alignment, Align(8));
+  } else if (GVSize > 2) {
+    // We might want to use a b32 load/store
+    Alignment = max(Alignment, Align(4));
+  } else if (GVSize > 1) {
+    // We might want to use a b16 load/store
+    Alignment = max(Alignment, Align(2));
+  }
+
+  return Alignment;
+}
+
+// Collect all kernel functions within which `GV` is used.
+static std::set<const Function *> collectUsedKernels(const GlobalVariable *GV) {
+  std::set<const Function *> UsedKernels;
+  SmallVector<const User *, 8> UserStack(GV->users());
+  SmallPtrSet<const User *, 8> VisitedUsers;
+
+  while (!UserStack.empty()) {
+    const auto *UU = UserStack.pop_back_val();
+
+    if (!VisitedUsers.insert(UU).second)
+      continue;
+
+    if (isa<GlobalVariable>(UU))
+      continue;
+
+    if (isa<Constant>(UU)) {
+      append_range(UserStack, UU->users());
+      continue;
+    }
+
+    const auto *I = dyn_cast<Instruction>(UU);
+    assert(I && "Expected an instruction");
+    const Function *F = I->getFunction();
+    if (AMDGPU::isModuleEntryFunctionCC(F->getCallingConv()))
+      UsedKernels.insert(F);
+  }
+
+  return UsedKernels;
+}
+
+static std::vector<const GlobalVariable *>
+getSortedUsedLDSGlobals(const Function *K) {
+  const Module *M = K->getParent();
+  const DataLayout &DL = M->getDataLayout();
+
+  // Collect all stastic LDS globals defined within the module.
+  std::vector<const GlobalVariable *> LDSGlobals =
+      AMDGPU::collectStaticLDSGlobals(M);
+
+  // Collect all those stastic LDS globals which are *used* within kernel K.
+  std::vector<const GlobalVariable *> UsedLDSGlobals;
+  for (const auto *GV : LDSGlobals) {
+    if (llvm::is_contained(collectUsedKernels(GV), K))
+      UsedLDSGlobals.push_back(GV);
+  }
+
+  // Sort LDS globals (which are used within kernel K) by alignment, descending,
+  // on ties, by size, descending, on ties, by name, lexicographical.
+  llvm::stable_sort(
+      UsedLDSGlobals,
+      [&](const GlobalVariable *LHS, const GlobalVariable *RHS) -> bool {
+        Align ALHS = fixAlignment(DL, LHS, AMDGPU::getAlign(DL, LHS));
+        Align ARHS = fixAlignment(DL, RHS, AMDGPU::getAlign(DL, RHS));
+        if (ALHS != ARHS) {
+          return ALHS > ARHS;
+        }
+
+        TypeSize SLHS = DL.getTypeAllocSize(LHS->getValueType());
+        TypeSize SRHS = DL.getTypeAllocSize(RHS->getValueType());
+        if (SLHS != SRHS) {
+          return SLHS > SRHS;
+        }
+
+        return LHS->getName() < RHS->getName();
+      });
+
+  // Module LDS which is possibly created by the "Lower Module LDS" pass, should
+  // be allocated at address 0, irrespective of its size and alignment.
+  std::vector<const GlobalVariable *> SortedLDSGlobals;
+  GlobalVariable *ModuleLDS = M->getGlobalVariable("llvm.amdgcn.module.lds");
+  if (ModuleLDS)
+    SortedLDSGlobals.push_back(ModuleLDS);
+
+  for (const auto *GV : UsedLDSGlobals) {
+    if (GV != ModuleLDS)
+      SortedLDSGlobals.push_back(GV);
+  }
+
+  return SortedLDSGlobals;
+}
+
 AMDGPUMachineFunction::AMDGPUMachineFunction(const MachineFunction &MF)
     : MachineFunctionInfo(), Mode(MF.getFunction()),
       IsEntryFunction(
           AMDGPU::isEntryFunctionCC(MF.getFunction().getCallingConv())),
       IsModuleEntryFunction(
           AMDGPU::isModuleEntryFunctionCC(MF.getFunction().getCallingConv())),
       NoSignedZerosFPMath(MF.getTarget().Options.NoSignedZerosFPMath) {
   const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(MF);
 
   // FIXME: Should initialize KernArgSize based on ExplicitKernelArgOffset,
   // except reserved size is not correctly aligned.
   const Function &F = MF.getFunction();
 
   Attribute MemBoundAttr = F.getFnAttribute("amdgpu-memory-bound");
   MemoryBound = MemBoundAttr.getValueAsBool();
 
   Attribute WaveLimitAttr = F.getFnAttribute("amdgpu-wave-limiter");
   WaveLimiter = WaveLimitAttr.getValueAsBool();
 
   CallingConv::ID CC = F.getCallingConv();
   if (CC == CallingConv::AMDGPU_KERNEL || CC == CallingConv::SPIR_KERNEL)
     ExplicitKernArgSize = ST.getExplicitKernArgSize(F, MaxKernArgAlign);
+
+  if (IsModuleEntryFunction) {
+    std::vector<const GlobalVariable *> SortedLDSGlobals =
+        getSortedUsedLDSGlobals(&F);
+
+    for (const auto *GV : SortedLDSGlobals)
+      allocateLDSGlobal(F.getParent()->getDataLayout(), *GV);
+  }
 }
 
 unsigned AMDGPUMachineFunction::allocateLDSGlobal(const DataLayout &DL,
                                                   const GlobalVariable &GV) {
   auto Entry = LocalMemoryObjects.insert(std::make_pair(&GV, 0));
   if (!Entry.second)
     return Entry.first->second;
 
   Align Alignment =
       DL.getValueOrABITypeAlignment(GV.getAlign(), GV.getValueType());
 
+  // Make sure that natural alignment of `GV` is correctly fixed.
+  Alignment = fixAlignment(DL, &GV, Alignment);
+
   /// TODO: We should sort these to minimize wasted space due to alignment
   /// padding. Currently the padding is decided by the first encountered use
   /// during lowering.
   unsigned Offset = StaticLDSSize = alignTo(StaticLDSSize, Alignment);
 
   Entry.first->second = Offset;
   StaticLDSSize += DL.getTypeAllocSize(GV.getValueType());
 
   // Update the LDS size considering the padding to align the dynamic shared

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

 bool userRequiresLowering(const SmallPtrSetImpl<GlobalValue *> &UsedList,
                           User *InitialUser);
 
 std::vector<GlobalVariable *>
 findVariablesToLower(Module &M, const SmallPtrSetImpl<GlobalValue *> &UsedList);
 
 SmallPtrSet<GlobalValue *, 32> getUsedList(Module &M);
 
+std::vector<const GlobalVariable *> collectStaticLDSGlobals(const Module *M);
+
 } // end namespace AMDGPU
 
 } // end namespace llvm
 
 #endif // LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H

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

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.
   SmallPtrSet<User *, 8> Visited;
   SmallVector<User *, 16> Stack;
   Stack.push_back(InitialUser);
 
   while (!Stack.empty()) {
     User *V = Stack.pop_back_val();
     Visited.insert(V);
 
bool userRequiresLowering(const SmallPtrSetImpl<GlobalValue *> &UsedList,
 
   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
SmallPtrSet<GlobalValue *, 32> getUsedList(Module &M) {
 
   TmpVec.clear();
   collectUsedGlobalVariables(M, TmpVec, false);
   UsedList.insert(TmpVec.begin(), TmpVec.end());
 
   return UsedList;
 }
 
+std::vector<const GlobalVariable *> collectStaticLDSGlobals(const Module *M) {
+  std::vector<const GlobalVariable *> StaticLDSGlobals;
+
+  for (const auto &GV : M->globals()) {
+    if (GV.getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {
+      // Skip non LDS global.
+      continue;
+    }
+
+    if (GV.hasExternalLinkage()) {
+      // HIP uses an unsized array `extern __shared__ T s[]` or similar
+      // zero-sized type in other languages to declare the dynamic shared
+      // memory variable whose size is not known at the compile time. They will
+      // be allocated by the runtime and placed directly after the statically
+      // allocated ones. Skip such dynamic shared memory variables.
+      assert(GV.getType()->isEmptyTy() &&
+             "Expected dynamic shared memory variable type to be empty");
+      continue;
+    }
+
+    if (!GV.hasInitializer()) {
+      // Static LDS globals should have "UndefValue" as an initializer. Skip LDS
+      // globals which do not have an initializer. ISEL pass will catch such
+      // static LDS globals and report error accordingly.
+      continue;
+    }
+
+    if (!isa<UndefValue>(GV.getInitializer())) {
+      // Static LDS globals should have "UndefValue" as an initializer. Skip LDS
+      // globals which have an initialzer but it is not "UndefValue". ISEL pass
+      // will catch such  static LDS globals and report error accordingly.
+      continue;
+    }
+
+    if (GV.isConstant()) {
+      // A constant LDS global cannot be allocated. Skip it.
+      continue;
+    }
+
+    StaticLDSGlobals.push_back(&GV);
+  }
+
+  return StaticLDSGlobals;
+}
+
 } // end namespace AMDGPU
 
 } // end namespace llvm

llvm/test/CodeGen/AMDGPU/GlobalISel/lds-global-value.ll

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=bonaire -verify-machineinstrs < %s | FileCheck %s
 ; TODO: Replace with existing DAG tests
 
 @lds_512_4 = internal unnamed_addr addrspace(3) global [128 x i32] undef, align 4
 @lds_4_8 = addrspace(3) global i32 undef, align 8
 
 define amdgpu_kernel void @use_lds_globals(i32 addrspace(1)* %out, i32 addrspace(3)* %in) #0 {
 ; CHECK-LABEL: use_lds_globals:
 ; CHECK:       ; %bb.0: ; %entry
 ; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
 ; CHECK-NEXT:    v_mov_b32_e32 v0, 4
 ; CHECK-NEXT:    s_mov_b32 m0, -1
-; CHECK-NEXT:    ds_read_b32 v3, v0 offset:4
+; CHECK-NEXT:    ds_read_b32 v3, v0
 ; CHECK-NEXT:    v_mov_b32_e32 v2, 9
 ; CHECK-NEXT:    s_waitcnt lgkmcnt(0)
 ; CHECK-NEXT:    s_add_u32 s0, s0, 4
 ; CHECK-NEXT:    s_addc_u32 s1, s1, 0
 ; CHECK-NEXT:    v_mov_b32_e32 v0, s0
 ; CHECK-NEXT:    v_mov_b32_e32 v1, s1
 ; CHECK-NEXT:    flat_store_dword v[0:1], v3
-; CHECK-NEXT:    v_mov_b32_e32 v0, 0
+; CHECK-NEXT:    v_mov_b32_e32 v0, 0x200
 ; CHECK-NEXT:    ds_write_b32 v0, v2
 ; CHECK-NEXT:    s_endpgm
 entry:
   %tmp0 = getelementptr [128 x i32], [128 x i32] addrspace(3)* @lds_512_4, i32 0, i32 1
   %tmp1 = load i32, i32 addrspace(3)* %tmp0
   %tmp2 = getelementptr i32, i32 addrspace(1)* %out, i32 1
   store i32 %tmp1, i32 addrspace(1)* %tmp2
   store i32 9, i32 addrspace(3)* @lds_4_8
   ret void
 }
 
 attributes #0 = { nounwind }

llvm/test/CodeGen/AMDGPU/lds-alignment.ll

 define amdgpu_kernel void @test_high_align_lds_arg(i8 addrspace(1)* %out, i8 addrspace(1)* %in, i8 addrspace(3)* align 64 %lds.arg) #1 {
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 64 %lds.arg, i8 addrspace(1)* align 64 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 64 %out, i8 addrspace(3)* align 64 %lds.arg, i32 38, i1 false)
   ret void
 }
 
 ; (7 * 8) + (39 * 4) = 212
 ; HSA-LABEL: {{^}}test_missing_alignment_size_2_order0:
-; HSA: workgroup_group_segment_byte_size = 212
+; HSA: workgroup_group_segment_byte_size = 216
 ; HSA: group_segment_alignment = 4
 define amdgpu_kernel void @test_missing_alignment_size_2_order0(i8 addrspace(1)* %out, i8 addrspace(1)* %in) #1 {
   %lds.missing.align.0.bc = bitcast [39 x i32] addrspace(3)* @lds.missing.align.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 4 %lds.missing.align.0.bc, i8 addrspace(1)* align 4 %in, i32 160, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 4 %out, i8 addrspace(3)* align 4 %lds.missing.align.0.bc, i32 160, i1 false)
 
   %lds.missing.align.1.bc = bitcast [7 x i64] addrspace(3)* @lds.missing.align.1 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.missing.align.1.bc, i8 addrspace(1)* align 8 %in, i32 56, i1 false)
define amdgpu_kernel void @test_round_size_3_order1(i8 addrspace(1)* %out, i8 addrspace(1)* %in) #1 {
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align16.0.bc, i32 38, i1 false)
 
   ret void
 }
 
 ; align 16, 32, 8
 ; 38 + (26 pad) + 38 + (10 pad) + 38 = 150
 ; HSA-LABEL: {{^}}test_round_size_3_order2:
-; HSA: workgroup_group_segment_byte_size = 150
+; HSA: workgroup_group_segment_byte_size = 134
 ; HSA: group_segment_alignment = 4
 define amdgpu_kernel void @test_round_size_3_order2(i8 addrspace(1)* %out, i8 addrspace(1)* %in) #1 {
   %lds.align16.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align16.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align16.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align16.0.bc, i32 38, i1 false)
 
   %lds.align32.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align32.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align32.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align32.0.bc, i32 38, i1 false)
 
   %lds.align8.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align8.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align8.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align8.0.bc, i32 38, i1 false)
 
   ret void
 }
 
 ; align 16, 8, 32
 ; 38 + (2 pad) + 38 + (2 pad) + 38
 ; HSA-LABEL: {{^}}test_round_size_3_order3:
-; HSA: workgroup_group_segment_byte_size = 118
+; HSA: workgroup_group_segment_byte_size = 134
 ; HSA: group_segment_alignment = 4
 define amdgpu_kernel void @test_round_size_3_order3(i8 addrspace(1)* %out, i8 addrspace(1)* %in) #1 {
   %lds.align16.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align16.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align16.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align16.0.bc, i32 38, i1 false)
 
   %lds.align8.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align8.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align8.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align8.0.bc, i32 38, i1 false)
 
   %lds.align32.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align32.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align32.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align32.0.bc, i32 38, i1 false)
 
   ret void
 }
 
 ; align 8, 32, 16
 ; 38 + (26 pad) + 38 + (2 pad) + 38 = 142
 ; HSA-LABEL: {{^}}test_round_size_3_order4:
-; HSA: workgroup_group_segment_byte_size = 142
+; HSA: workgroup_group_segment_byte_size = 134
 ; HSA: group_segment_alignment = 4
 define amdgpu_kernel void @test_round_size_3_order4(i8 addrspace(1)* %out, i8 addrspace(1)* %in) #1 {
   %lds.align8.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align8.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align8.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align8.0.bc, i32 38, i1 false)
 
   %lds.align32.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align32.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align32.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align32.0.bc, i32 38, i1 false)
 
   %lds.align16.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align16.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align16.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align16.0.bc, i32 38, i1 false)
 
   ret void
 }
 
 ; align 8, 16, 32
 ; 38 + (10 pad) + 38 + (2 pad) + 38 = 126
 ; HSA-LABEL: {{^}}test_round_size_3_order5:
-; HSA: workgroup_group_segment_byte_size = 126
+; HSA: workgroup_group_segment_byte_size = 134
 ; HSA: group_segment_alignment = 4
 define amdgpu_kernel void @test_round_size_3_order5(i8 addrspace(1)* %out, i8 addrspace(1)* %in) #1 {
   %lds.align8.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align8.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align8.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)
   call void @llvm.memcpy.p1i8.p3i8.i32(i8 addrspace(1)* align 8 %out, i8 addrspace(3)* align 8 %lds.align8.0.bc, i32 38, i1 false)
 
   %lds.align16.0.bc = bitcast [38 x i8] addrspace(3)* @lds.align16.0 to i8 addrspace(3)*
   call void @llvm.memcpy.p3i8.p1i8.i32(i8 addrspace(3)* align 8 %lds.align16.0.bc, i8 addrspace(1)* align 8 %in, i32 38, i1 false)

llvm/test/CodeGen/AMDGPU/local-memory.amdgcn.ll

 ; GCN: .long 47180
 ; GCN-NEXT: .long 32900
 
 ; GCN-LABEL: {{^}}local_memory_two_objects:
 ; GCN: v_lshlrev_b32_e32 [[ADDRW:v[0-9]+]], 2, v0
 ; CI-DAG: v_sub_i32_e32 [[SUB:v[0-9]+]], vcc, 0, [[ADDRW]]
 ; CI-DAG: ds_write2_b32 [[ADDRW]], {{v[0-9]+}}, {{v[0-9]+}} offset1:4
 ; SI-DAG: ds_write2_b32 [[ADDRW]], {{v[0-9]+}}, {{v[0-9]+}} offset1:4
-; SI-DAG: v_sub_i32_e32 [[SUB0:v[0-9]+]], vcc, 28, [[ADDRW]]
+; SI-DAG: v_sub_i32_e32 [[SUB0:v[0-9]+]], vcc, 12, [[ADDRW]]
 
 ; GCN: s_barrier
 
-; SI-DAG: v_sub_i32_e32 [[SUB1:v[0-9]+]], vcc, 12, [[ADDRW]]
+; SI-DAG: v_sub_i32_e32 [[SUB1:v[0-9]+]], vcc, 28, [[ADDRW]]
 ; SI-DAG: ds_read_b32 v{{[0-9]+}}, [[SUB0]]
 ; SI-DAG: ds_read_b32 v{{[0-9]+}}, [[SUB1]]
 ; CI: ds_read2_b32 {{v\[[0-9]+:[0-9]+\]}}, [[SUB]] offset0:3 offset1:7
 
 define amdgpu_kernel void @local_memory_two_objects(i32 addrspace(1)* %out) #0 {
 entry:
   %x.i = call i32 @llvm.amdgcn.workitem.id.x()
   %arrayidx = getelementptr inbounds [4 x i32], [4 x i32] addrspace(3)* @local_memory_two_objects.local_mem0, i32 0, i32 %x.i

llvm/test/CodeGen/AMDGPU/promote-alloca-padding-size-estimate.ll

 
 
 @lds0 = internal unnamed_addr addrspace(3) global [32 x <4 x i32>] undef, align 16
 @lds2 = internal unnamed_addr addrspace(3) global [32 x i64] undef, align 8
 @lds1 = internal unnamed_addr addrspace(3) global [73 x i32] undef, align 4
 
 
 ; GCN-LABEL: {{^}}promote_alloca_size_order_0:
-; GCN: workgroup_group_segment_byte_size = 1060
+; GCN: workgroup_group_segment_byte_size = 1072
 define amdgpu_kernel void @promote_alloca_size_order_0(i32 addrspace(1)* nocapture %out, i32 addrspace(1)* nocapture %in, i32 %idx) #0 {
 entry:
   %stack = alloca [5 x i32], align 4, addrspace(5)
   %tmp0 = load i32, i32 addrspace(1)* %in, align 4
   %arrayidx1 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %tmp0
   store i32 4, i32 addrspace(5)* %arrayidx1, align 4
   %arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %in, i32 1
   %tmp1 = load i32, i32 addrspace(1)* %arrayidx2, align 4