[AMDGPU] Replace uses of LDS globals within non-kernel functions by pointers.

[hsmhsm]

Summary

One of the memory types being supported within AMD GPU memory hierarchy is
shared memory, also called Local Data Share or LDS for short. LDS memory
is the second fastest memory in the AMD GPU memory hierarchy (with register
file being fastest available memory in the hierarchy). Being faster also
means LDS memory is comparatively costlier and hence is a limited available
memory resource.

Being global scoped, an LDS variable is accessible within kernel functions
and non-kernel functions, but two different kernel execution paths, say
called from two kernels K1 and K2, cannot access the same instance of an LDS
variable, say L. Both K1 and K2 has to own its own instance of L. This puts
some challenges, especially to lower the LDS variables used within non-kernel
functions.

So, the pass - "Lower Module LDS" lowers the LDS globals by packing them
within in a struct type, and by creating an instance of that struct type
within every kerenl at address zero. Though, the pass - "Lower Module LDS"
makes some effort to minimize unnecessary LDS allocation, it is limited by
means of the fundamental basis and assumption upon which the pass is
implemented.

The current pass acts as an helping aid to the pass - "Lower Module LDS" with
the intention of minimizing unnecessary LDS allocation as much as possible.

The main idea behind the current pass is:

(1) To identify the LDS globals used within non-kernel function scope and
global scope.
(2) To push the use of all the above identified LDS globals to kernel
function scope by initializing their addresses to newly created LDS
global pointer variables (within kernel functions).
(3) To replace the uses of original LDS globals within non-kernel functions
by their pointer counter-parts.
(4) This way, the transformation makes sure that the pass "Lower Module LDS"
packs only pointer variables within struct type, and hence significantly
minimizes unnecessary LDS allocation, espacically when the original LDS
globals are big arrays (as this is the common LDS use case).

NOTE: The pass - "Lower Module LDS" now has a tight dependency on the current pass, and the current pass should always be run before running the pass "Lower Module LDS". Running the pass "Lower Module LDS" alone may lead to surprizing results.

[hsmhsm]

[0] Started re-implementing from scratch again.
[1] Added a new pass, namely, amdgpu-lower-function-local-lds.
[2] Implemented required initial plumbing work for both old and new pass

managers.

[3] An option, namely, amdgpu-enable-function-local-lds-lowering is

added, when passed, it enables the pass.

[hsmhsm]

Started to implement the feature from scratch again. The previous experience tells me that - "a single very big patch is very problamatic and confusing for a meanigful review process". Hence this time, I am planning to submit small patches (time to time) which can be reasonably reviewed. This first patch implements the following.

[1] Add new pass, namely, amdgpu-lower-function-local-lds.
[2] Implement required initial plumbing work for both old and new pass managers.
[3] Add an option, namely, amdgpu-enable-function-local-lds-lowering, when passed, it enables the pass.

[arsenm]

You can just use the done checkbox, you don't need to comment on each point

llvm/lib/Target/AMDGPU/AMDGPUDeviceScopeSharedVariable.cpp

1510 ↗

(On Diff #313659)

You should only try to preserve things that are important, otherwise you are adding cost and complexity for no benefit

1512 ↗

(On Diff #313659)

The IR is language independent and none of the constructs here are tied to a language

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h

38

This isn't a user exposed flag, and there shouldn't be a need for users to set one.

[hsmhsm]

Based on the Matt's comment for previous patch, some changes are done w.r.t the handling of the guarding flag - amdgpu-enable-function-local-lds-lowering.

Though Matt is against the usage of any guarding flag for this pass, I personally feel the need of it for below two reasons.

(1) Presence of this pass means, we should disable the forcefull inlining as it is done within the pass - amdgpu-always-inline. Otherwise, this pass does not make any sense at all. It is a better idea to disable this forcefull inlining via a flag.

(2) In case of any emergency issue within this pass, customer should have an handy approach to disable the pass and temporarily move on until the fix available.

So the amdgpu-enable-function-local-lds-lowering is a hidden flag, and it is enabled by default. It works as below:

(1) Default behavoir is to run the pass as shown below.

• Old pass manager:

mahesha@brego:[tmp]$ hipcc main.cpp
Running the pass - LowerFunctionLocalLDS
mahesha@brego:[tmp]$

• New pass manager:

mahesha@brego:[tmp]$ hipcc -fexperimental-new-pass-manager main.cpp
Running the pass - LowerFunctionLocalLDS
mahesha@brego:[tmp]$

(2) The pass will not run when it is explicitly turned off as shown below.

• Old pass manager:

mahesha@brego:[tmp]$ hipcc -mllvm --amdgpu-enable-function-local-lds-lowering=false main.cpp
mahesha@brego:[tmp]$

• New pass manager:

mahesha@brego:[tmp]$ hipcc -fexperimental-new-pass-manager  -mllvm --amdgpu-enable-function-local-lds-lowering=false main.cpp
mahesha@brego:[tmp]$

[hsmhsm]

Build all the required data structures which will be later used to lower function local LDS. Below are the data structures being built.

[1] Kernel Set - Holds all the kernels in the module
[2] Function Local LDS Set - Holds all the function local LDS from all functions
[3] Function Address Taken Set - Holds all the functions whose address is taken within the module

[4] LDS to Function Map - Maps each function local LDS to a function within which the LDS is defined
[5] Function to LDS Map - Reverse of above map, which maps each functon F to a SET of LDS which are defined within F

[6] Kernel to Callee Map - Maps each kernel K to a SET of functions which define LDS and there exists call graph path from K to these functions.
[7] Kernel to LDS - Maps each kernel K to a set of function local LDS which are supposed to be lowered w.r.t K.

Data structures [1], [2], and [3] are built by iterating over the globals and functions defined within the module.
Data structures [4] and [5] are built using BOTTOM-UP based on the use list of function local LDS.
Data structure [6] is built using TOP-DOWN via call graph traversal.
Data structure [7] is built using the result of above BOTTOM-UP and TOP-DOWN constructed data structures.

[hsmhsm]

Add missing "static" keyword to a function isKernel().

[hsmhsm]

Added a FIXME comment.

[hsmhsm]

Add missing explicit keyword for constructor.

[hsmhsm]

Fix few spell mistakes in comments.

[hsmhsm]

Fix few spell mistakes in comments.

[hsmhsm]

Fix comments.

[hsmhsm]

Re-arrange code for more readability.

[hsmhsm]

Fixed clang-tidy warnings.

[hsmhsm]

Code re-organization.

[hsmhsm]

Corrected few comments.

[hsmhsm]

Fixed one of the FIXME comments which is associated with indirect calls.

[hsmhsm]

Make use of llvm append_range() api.

[hsmhsm]

Improvements to code at few places.

[arsenm]

All tests are now missing

llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp

32 ↗

(On Diff #319026)

This function is pointless, just directly use isModuleEntryFunctionCC

41–44 ↗

(On Diff #319026)

This function is useless. Assert strings also don't need to end in \n

56 ↗

(On Diff #319026)

cast<>, don't dyn_cast and assert

66 ↗

(On Diff #319026)

Pointless comment

68 ↗

(On Diff #319026)

Pointless comment

74–75 ↗

(On Diff #319026)

Pointless comment

82 ↗

(On Diff #319026)

Extra private

92 ↗

(On Diff #319026)

Extra private

93 ↗

(On Diff #319026)

Typo unhanlded

111 ↗

(On Diff #319026)

.contains

142 ↗

(On Diff #319026)

Extra private

233 ↗

(On Diff #319026)

Copy of set unnecessary

320 ↗

(On Diff #319026)

Don't need all these newlines in assert strings

376–379 ↗

(On Diff #319026)

I think you're overcomplicating the CallGraph usage by ignoring most of what it gives you. You should be able to just iterate directly through the CallGraph to get functions reachable from the parent

410–411 ↗

(On Diff #319026)

This concept doesn't quite work for the IR. The same global can appear in multiple functions

449 ↗

(On Diff #319026)

isa<>, no \n

479 ↗

(On Diff #319026)

Should not be checking the function name. Should just skip all declarations

492–495 ↗

(On Diff #319026)

Return !Kernels.empty()

502 ↗

(On Diff #319026)

Probably should skip declarations. Also not sure about the linkage check

506–509 ↗

(On Diff #319026)

Return !empty()

555–558 ↗

(On Diff #319026)

Return !empty()

561–571 ↗

(On Diff #319026)

Don't understand the point of this stub function

[hsmhsm]

Fixed review comments (by Matt).

[hsmhsm]

llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp

111 ↗

(On Diff #319026)

The data structures, ValueMap<>, SmallPtrSet<> do not have member function - .contains(). W.r.t std::set<>, this member function is supported in C++20.

376–379 ↗

(On Diff #319026)

As far as I understand it, llvm CallGraph infrastructure does not provide any facility as such. Implementer needs to explicitly iterate the callees of the caller.

410–411 ↗

(On Diff #319026)

My understanding is that - scope of the shared variable is function/statement block scope. It is not available to access outside this scope. It is just that we implement it as global, just like how the local static variables are implemented in C/C++?

Can you give an example of the use-case that you are claiming?

502 ↗

(On Diff #319026)

The linkage test is required to ignore the dynamic shared variables like the one defined as extern __shared__ int dy_sm[]; where size of dy_sm is not known at compile time, but is passed as one of the kernel execution configuration parameters at run time.

561–571 ↗

(On Diff #319026)

This is a driver function, it looks like a stub now, since implementation is not complete yet. Once this patch is accepted, next step is to (1) define kernel specific LDS layouts (2) create 2D offset table and (3) add new implicit argument.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp

125–128 ↗

(On Diff #319316)

Should just inline this function

182 ↗

(On Diff #319316)

Should just inline this function

111 ↗

(On Diff #319026)

There is an llvm::is_contained. Also why use std::set? You randomly switch set types around here

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp

168 ↗

(On Diff #319316)

Should just inline this function

205 ↗

(On Diff #319316)

Copy here, just directly use this in the for loop

410–411 ↗

(On Diff #319026)

The IR has absolutely no concept of these scopes. The global variables have global scope and no restriction on where their uses can appear. Whether or not this directly corresponds to a direct language feature is unimportant. Some IPO transforms can push global variable references into other functions.

The example is just two functions that refer to the same variable:

@lds = ...

define void @func0() {
  store i32 0, i32* @lds
  ret void
}

define void @func1() {
  store i32 0, i32* @lds
  ret void
}

502 ↗

(On Diff #319026)

That's a function of it having 0 size, not the linkage

561–571 ↗

(On Diff #319026)

If it's going to be split, I'd rather see the full stack for the review

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp

376–379 ↗

(On Diff #319026)

The CallGraph as a whole gives you the functions reachable from each other. I don't think you need to do a stack walk to find the callees. You don't need to care about which functions specifically call which, just that they are all connected

[hsmhsm]

Pushing renamed code (not for review).

[hsmhsm]

Fixed Matt's comments.

[hsmhsm]

Save current work.

[hsmhsm]

Save current work.

[jdoerfert]

Could we get some tests *and* a commit message that explains what this is supposed to do.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPULowerLDSGlobal.cpp

312

You should only need to do the use replacement, you aren't changing the types of the instructions so cloning/hacking on them shouldn't be needed

447–449

Don't need this, the IR would have failed the verifier to get here

468–469

Should not const_cast

518–519

Too much auto for me

554

Should use lowercase, period separator naming convention with an llvm.amdgcn prefix

566

Should use lowercase, period separator naming convention with an llvm.amdgcn prefix

609

This still needs to add in alignment padding

669

Braces, Also can use range loop

743

I don't see why you need to build your own stack. The call graph already found the reachable functions for you

777

I think trying to handle callees is left for a later patch. Additionally, I think this should be the CallGraph analysis's responsibility to deal with

855–858

The callgraph should already give this to you. Iterating the call graph should give you all of the functions you care about. You don't actually need to worry about which functions call which, since you need to touch every function in the SCC

987–988

I would swap the order of these checks

[hsmhsm]

In D91516#2560020, @jdoerfert wrote:

Could we get some tests *and* a commit message that explains what this is supposed to do.

This is WIP, will add test and commit messages at the end before final review.

[hsmhsm]

Address Matt's comments.

[hsmhsm]

llvm/lib/Target/AMDGPU/AMDGPULowerLDSGlobal.cpp

312

Added FIXME comment, will see how to fix it.

609

Added FIXME comment, will see how to fix it.

743

Added FIXME comment, will see how to fix it.

777

Added FIXME comment, will see how to fix it.

855–858

Added FIXME comment, will see how to fix it.

[hsmhsm]

Assert that both caller and callee appear in same translation unit.

[hsmhsm]

Address clang-tidy warnings.

[hsmhsm]

Remove over created auto variables.

[jdoerfert]

In D91516#2561383, @hsmhsm wrote:

In D91516#2560020, @jdoerfert wrote:

Could we get some tests *and* a commit message that explains what this is supposed to do.

This is WIP, will add test and commit messages at the end before final review.

I generally would recommend against that but I guess you can use phabricator this way.
However, once people can figure out what this is actually supposed to do, they might effectively restart the entire review process if the design is questioned.
I say this because I have a hunch, or maybe a hope, about the intent of this patch. If it would be that, I'd very much like this to be a generic, non-AMDGPU pass. I might be wrong about what this does, and that is what I'd like to figure out rather sooner than later.

[hsmhsm]

In D91516#2562027, @jdoerfert wrote:

In D91516#2561383, @hsmhsm wrote:

In D91516#2560020, @jdoerfert wrote:

Could we get some tests *and* a commit message that explains what this is supposed to do.

This is WIP, will add test and commit messages at the end before final review.

I generally would recommend against that but I guess you can use phabricator this way.
However, once people can figure out what this is actually supposed to do, they might effectively restart the entire review process if the design is questioned.
I say this because I have a hunch, or maybe a hope, about the intent of this patch. If it would be that, I'd very much like this to be a generic, non-AMDGPU pass. I might be wrong about what this does, and that is what I'd like to figure out rather sooner than later.

OK let's wait for the complete patch then. I will only going to push complete patch next time. But, it may take some time since there are some major hurdles to overcome.

[hsmhsm]

Implemented a new approach based on initializing LDS globals to pointers.

[JonChesterfield]

This is much more complicated than I expected. Is the large amount of comments largely from a previous patch doing different things that has been hammered into this one?

@jdoerfert the transform I think this is intended to do is:

• find a large shared variable used from a function
• add a new void*, also in shared, pointing to it
• initialize that void* only in kernels that can call functions that use the large variable
• replace all uses with

That means, on amdgcn, the large variable only costs LDS space in kernels that definitely use it. I don't know how cuda lowers shared accesses from functions, it could plausibly benefit from the same transform.

[JonChesterfield]

I can't work out which LDS variables you intend to replace with pointers from the code. Could you spell out what the condition under which you intend to replace one is?

llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp

177 ↗

(On Diff #332541)

Why not isa<GlobalVariable> / function needs a different name

204 ↗

(On Diff #332541)

Functions define hasAddressTaken, but also I don't think this pass needs to distinguish between direct and indirect calls

224 ↗

(On Diff #332541)

This I haven't read yet, but it looks like far too much state. Expected a set of LDS globals called 'toReplaceWithPointer' or similar instead of all the maps

384 ↗

(On Diff #332541)

Why do we want to replace constexpr with instructions? This comment contradicts the implementation

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

20 ↗

(On Diff #332541)

Perhaps name the new files after LDS to make it clearer that they're used for LDS lowering an optimisation, not necessarily general purpose.

Also move the functions out in a separate commit, without changes to their implementation, as that improves the signal/noise of the functional change.

59 ↗

(On Diff #332541)

e.g. I recognise this as newly introduced by the comment, but in phab it's hard to distinguish from things that haven't changed

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

16 ↗

(On Diff #332541)

Include list should be limited to those that are used by the header, with the ones used by the source included there

[hsmhsm]

In D91516#2643733, @JonChesterfield wrote:

This is much more complicated than I expected.

We need to *really* discuss, what is complicated here and what is violated here from the internal email discussions.

Is the large amount of comments largely from a previous patch doing different things that has been hammered into this one?

No, nothing is hammered from the previous patch. The current patch is implementing what is planned via internal emial discussion.

@jdoerfert the transform I think this is intended to do is:

• find a large shared variable used from a function
• add a new void*, also in shared, pointing to it
• initialize that void* only in kernels that can call functions that use the large variable
• replace all uses with

No, the intended implementation plan which is implemented here is as follows.

(1) Identify the LDS globals (whether large or small) which are used within non-kernel function scope and in global scope.
(2) Create new LDS glboals of i16 type corresponding to every LDS global identified above. The i16 typed LDS globals act as pointers to corresponding original LDS globals.
(3) push the *use* of above identified LDS globals to kernels by adding instructions within the kernels which initialize the address of original LDS globals to their respective pointers. This is will make sure that per kernel LDS allocation for these LDS globals correctly happen.
(4) Within non-kernel functions, replace the *use* of original LDS globals by thier respective pointers.
(5) Keep the global scope use of original LDS globals unchanged since now they should work automatically as the use of these original LDS globals (pointer initialiation) also there within all kernels and hence it will semantically work correct as expected because of per kernel LDS allocation for these LDS globals.

That means, on amdgcn, the large variable only costs LDS space in kernels that definitely use it. I don't know how cuda lowers shared accesses from functions, it could plausibly benefit from the same transform.

Let's not bother about how CUDA handles it since there is lot of differences here. And focus on only AMDGCN.

[hsmhsm]

In D91516#2644163, @JonChesterfield wrote:

I can't work out which LDS variables you intend to replace with pointers from the code. Could you spell out what the condition under which you intend to replace one is?

All those LDS globals which are used within non-kernel functions and within global scope requires pointer initialization within kernels.

[hsmhsm]

llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp

177 ↗

(On Diff #332541)

Because LDS would be nested within the const expr within global scope use.

204 ↗

(On Diff #332541)

I am not getting this comment, probably we can discuss it offline.

224 ↗

(On Diff #332541)

These maps are required for the logic where we really need to restrict the LDS set for kernel based on kernel excecution paths.

384 ↗

(On Diff #332541)

Again not clear about what you intended here - let's take it offline.

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

20 ↗

(On Diff #332541)

will think about it.

59 ↗

(On Diff #332541)

not sure what you mean here. Let's discuss offline.

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

16 ↗

(On Diff #332541)

agree.

[hsmhsm]

Hi Jon,

I have replied to some of your review comments, and few other comments require internal discussion for better and quick unambiguous conclusion. I am expecting a response from you.

[JonChesterfield]

In D91516#2655086, @hsmhsm wrote:

In D91516#2643733, @JonChesterfield wrote:

This is much more complicated than I expected.
Is the large amount of comments largely from a previous patch doing different things that has been hammered into this one?

No, nothing is hammered from the previous patch. The current patch is implementing what is planned via internal emial discussion.

This review dates from November 16 last year and contains hundreds of review comments against code that may or may not still be in the latest revision, this being diff #34 at time of commenting. If I'm following along successfully, the design has changed significantly and repeatedly during that process. It is therefore very difficult to determine what the design intent behind the current revision is. That is what I mean by 'previous patch has been hammered into this one'.

[JonChesterfield]

The algorithm I had in mind was along the lines of:

for each LDS variable:
  if should-transform
    create 16 bit integer in LDS
    initialize that global with (constexpr) address of variable
    replace all uses of variable with a (constexpr) access through new pointer

where
should-transform:
 if (sizeof) < 8ish return false
 if used by instruction in indirectly called function return false
 if only used by kernels return false
 probably other exclusions
 return true

llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp

384 ↗

(On Diff #332541)

I was thinking the introduced 16 bit pointers will be initialised with constexpr from the corresponding variable.

This patch presently initialises them with undef, which I think thwarts using constexpr everywhere, and means we insert stores in the kernel entry basic block here.

If we fix the back end to handle LDS variables with initializers (at least the simple case of only used from kernel and initialized with address of some other variable), then quite a lot of the complexity of this patch drops out.

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

195–202

I think this test reads better as proposed here - 'enable-lower-module-lds=true' is better than 'disable-lower-module-lds=false'. Separable from the rest of this patch, we could land a patch that just inverts that commandline flag and updates the tests to match. That removes some noise from this review.

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

22 ↗

(On Diff #332541)

This test should only check the behaviour of lower-module-lds. Separate tests check the behaviour of amdgpu-replace-lds-use-with-pointer.

Equally, running amdgpu-lower-module-lds by itself should not automatically run amdgpu-replace-lds-use-with-pointer and vice versa.

llvm/test/CodeGen/AMDGPU/replace_lds_report_error_no_func_def.ll

3 ↗

(On Diff #332541)

This is an error in the implementation, not something that should have a test checking the implementation is broken. Instead of assuming the definition of both are in the same module and crashing if they aren't, the pass should ignore a variable which doesn't meet that requirement.

llvm/test/CodeGen/AMDGPU/replace_lds_test_direct_call_misc.ll

15 ↗

(On Diff #332541)

These tests would be more robust if the new pointer was named based on the global it is intended to reference, as then the regex can check that we created load from the correct pointer (as opposed to just one of the new pointers).

[hsmhsm]

In D91516#2682974, @JonChesterfield wrote:

The algorithm I had in mind was along the lines of:

for each LDS variable:
  if should-transform
    create 16 bit integer in LDS
    initialize that global with (constexpr) address of variable
    replace all uses of variable with a (constexpr) access through new pointer

where
should-transform:
 if (sizeof) < 8ish return false
 if used by instruction in indirectly called function return false
 if only used by kernels return false
 probably other exclusions
 return true

I think, we all of us, who are involved in discussing about the functionalities related to this patch are not on the same page. First, we need to internally discuss it and make sure that we are all on the same page, before I start making any further changes to this patch.

[hsmhsm]

By the way, this patch has gone through too many revisions, and is becoming too complex to go back to any previous history of the patch when required. So, I think, it is better to abandon this patch, start on a clean slate with fresh new patch. If I do not get any objection to abandon this patch, then I will be abonding it in a day or two.

[hsmhsm]

As mentioned eariler, I am abandoning this patch. Let's start with a clean slate, and decide on the implementation.

llvm/lib/Target/AMDGPU/AMDGPU.h

 extern char &AMDGPUOpenCLEnqueuedBlockLoweringID;
 
 void initializeGCNRegBankReassignPass(PassRegistry &);
 extern char &GCNRegBankReassignID;
 
 void initializeGCNNSAReassignPass(PassRegistry &);
 extern char &GCNNSAReassignID;
 
+ModulePass *createAMDGPULowerLDSGlobalPass();
+void initializeAMDGPULowerLDSGlobalPass(PassRegistry &);
+extern char &AMDGPULowerLDSGlobalID;
+struct AMDGPULowerLDSGlobalPass : PassInfoMixin<AMDGPULowerLDSGlobalPass> {
+  AMDGPULowerLDSGlobalPass() {}
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &AM);
+};
+
 namespace AMDGPU {
 enum TargetIndex {
   TI_CONSTDATA_START,
   TI_SCRATCH_RSRC_DWORD0,
   TI_SCRATCH_RSRC_DWORD1,
   TI_SCRATCH_RSRC_DWORD2,
   TI_SCRATCH_RSRC_DWORD3
 };

llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp

static bool alwaysInlineImpl(Module &M, bool GlobalOpt) {
   // objects defined in functions. LDS is always allocated by a kernel, and it
   // is difficult to manage LDS usage if a function may be used by multiple
   // kernels.
   //
   // OpenCL doesn't allow declaring LDS in non-kernels, so in practice this
   // should only appear when IPO passes manages to move LDs defined in a kernel
   // into a single user function.
 
+  if (!AMDGPUTargetMachine::EnableLDSGlobalLowering) {
     for (GlobalVariable &GV : M.globals()) {
       // TODO: Region address
       unsigned AS = GV.getAddressSpace();
       if (AS != AMDGPUAS::LOCAL_ADDRESS && AS != AMDGPUAS::REGION_ADDRESS)
         continue;
 
       recursivelyVisitUsers(GV, FuncsToAlwaysInline);
     }
+  }
 
   if (!AMDGPUTargetMachine::EnableFunctionCalls || StressCalls) {
     auto IncompatAttr
       = StressCalls ? Attribute::AlwaysInline : Attribute::NoInline;
 
     for (Function &F : M) {
       if (!F.isDeclaration() && !F.use_empty() &&
           !F.hasFnAttribute(IncompatAttr)) {

llvm/lib/Target/AMDGPU/AMDGPULowerLDSGlobal.cpp

+//===-- AMDGPULowerLDSGlobal.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
+//
+//===----------------------------------------------------------------------===//
+//
+// TODO
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPU.h"
+#include "Utils/AMDGPUBaseInfo.h"
+#include "llvm/ADT/SetVector.h"
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/ADT/SmallSet.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/Analysis/CallGraph.h"
+#include "llvm/CodeGen/TargetPassConfig.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/MDBuilder.h"
+#include "llvm/IR/Module.h"
+#include "llvm/IR/ValueMap.h"
+#include "llvm/InitializePasses.h"
+#include "llvm/Transforms/Utils/Cloning.h"
+#include <map>
+#include <set>
+
+#define DEBUG_TYPE "amdgpu-lower-lds-global"
+
+using namespace llvm;
+
+// Helper function around `ValueMap` to detect if an element exists within it.
+template <typename R, typename E>
+static bool contains(R &&VMap, const E &Element) {
+  return VMap.find(Element) != VMap.end();
+}
+
+// Helper function to get the required alignment for the input LDS variable
+// `GV`.
+static Align getAlign(GlobalVariable *GV) {
+  return GV->getAlign().getValueOr(
+      GV->getParent()->getDataLayout().getPreferredAlign(GV));
+}
+
+namespace {
+
+class LowerLDSGlobalImpl {
+  Module &M;
+  LLVMContext &Ctx;
+  const DataLayout &DL;
+  CallGraph CG;
+  Twine PrefixStr;
+
+  // Holds kernels defined within the module `M`.
+  SmallPtrSet<Function *, 16> Kernels;
+
+  // Holds LDS globals defined within the module `M`.
+  SmallPtrSet<GlobalVariable *, 16> LDSGlobals;
+
+  // Holds call graph nodes associated with the functions whose addresses are
+  // taken within the module.
+  SmallPtrSet<CallGraphNode *, 16> AddressTakenSet;
+
+  // Associates LDS global to a list of functions which references that LDS.
+  ValueMap<GlobalVariable *, SmallPtrSet<Function *, 8>> LDSGlobalToAccessors;
+
+  // Associates function to a list of LDS globals which are referenced within
+  // that function.
+  ValueMap<Function *, SmallPtrSet<GlobalVariable *, 8>> AccessorToLDSGlobals;
+
+  // Associates kernel to a list of functions which are reachable from that
+  // kernel.
+  ValueMap<Function *, SmallPtrSet<Function *, 8>> KernelToCallees;
+
+  // Associates kernel to a list LDS globals which are referenced along the call
+  // graph paths associated with that kernel.
+  ValueMap<Function *, SmallPtrSet<GlobalVariable *, 8>> KernelToLDSGlobals;
+
+  // Associates kernel to a newly created per kernel LDS layout.
+  ValueMap<Function *, GlobalVariable *> KernelToLDSLayout;
+
+  // Associates kernel to a map which maps LDS globals to corresponding offsets.
+  ValueMap<Function *, std::map<GlobalVariable *, uint32_t>>
+      KernelToLDSToOffset;
+
+  // Associates kernel to instruction insertion point within that kernel.
+  ValueMap<Function *, Instruction *> KernelToInstInsertPt;
+
+  // Holds pointers to LDS layouts.
+  GlobalVariable *LDSLayouts;
+
+  // Holds 2D LDS offset table.
+  GlobalVariable *LDSOffsetTable;
+
+  // Associates kernel to unique integer id.
+  std::map<Function *, uint32_t> KernelToID;
+
+  // Associates an integer id uniquely to kernel.
+  std::map<GlobalVariable *, uint32_t> LDSToID;
+
+  // Associates LDS global to unique integer id.
+  std::map<uint32_t, Function *> IDToKernel;
+
+  // Associates an integer id uniquely to LDS global.
+  std::map<uint32_t, GlobalVariable *> IDToLDS;
+
+public:
+  explicit LowerLDSGlobalImpl(Module &M)
+      : M(M), Ctx(M.getContext()), DL(M.getDataLayout()), CG(CallGraph(M)),
+        PrefixStr(Twine("lds.lower.")) {
+    // Collect the functions whose address is taken within the module.
+    collectAddressTakenFunctions();
+  }
+
+  // Entry-point function.
+  bool lower();
+
+private:
+  //===--------------------------------------------------------------------===//
+  // Methods which aid in creating new instructions which assess LDS layouts.
+  //===--------------------------------------------------------------------===//
+
+  // Convert `ConstantExpr CE` to `Instruction`, and update users of `CE` to use
+  // `Instruction`.
+  Instruction *replaceConstExprByInst(ConstantExpr *CE);
+
+  // Within instruction `I` replace the use(s) of `LDS` by `LayoutAccessInst`.
+  void replaceInstWhichUsesLDS(Instruction *I, GlobalVariable *LDS,
+                               Value *LayoutAccessInst,
+                               SmallPtrSetImpl<Instruction *> &ToBeErasedInsts);
+
+  // Replace all the users of original LDS global `LDS` with new ones which
+  // access corresponding LDS layouts.
+  void replaceAllUsersOfLDS(
+      GlobalVariable *LDS,
+      ValueMap<Function *, Instruction *> &FunctionToLDSLayoutAccessInst);
+
+  // Add instructions to access LDS layouts corresponding to LDS global `LDS`.
+  bool accessLDSLayouts(
+      GlobalVariable *LDS,
+      ValueMap<Function *, Instruction *> &FunctionToLDSLayoutAccessInst);
+
+  // Add instructions within kernels to initialize LDS layouts array to hold the
+  // starting address of LDS layouts.
+  void initializeArrayOfLDSLayouts();
+
+  // Insert various required instructions within kernels and within callees
+  // which aid accessing of newly inserted LDS layouts.
+  void insertNewInstructions();
+
+  //===--------------------------------------------------------------------===//
+  // Methods which aid in creating new globals like kernel specific LDS layouts,
+  // 2D LDS offset table, etc.
+  //===--------------------------------------------------------------------===//
+
+  // Create 2D LDS offset table which will be referenced at run time to access
+  // the LDS specific offset within kernel specific LDS layout.
+  void constructLDSOffsetTable();
+
+  // Create 1D array which holds LDS layouts which will be referenced at run
+  // time to access kernel specific LDS layout.
+  void constructArrayOfLDSLayouts();
+
+  // Create a new LDS layout corresponding to kernel `K` whose size will be of
+  // `LayoutSize`.
+  GlobalVariable *constructLDSLayout(Function *K, uint32_t LayoutSize);
+
+  // Given a set of LDS globals, associate each LDS from the set to a
+  // corresponding offset within the kernel specific LDS layout.
+  uint32_t computeOffsets(SmallPtrSetImpl<GlobalVariable *> &LDSSet,
+                          std::map<GlobalVariable *, uint32_t> &LDSToOffSet);
+
+  // Create unified LDS layouts for each kernel.
+  void constructLDSLayouts();
+
+  // Insert various required globals (including per kernel specific unified LDS
+  // layouts) into the module.
+  void insertNewGlobals();
+
+  //===--------------------------------------------------------------------===//
+  // Methods which aid in preparing for lowering.
+  //===--------------------------------------------------------------------===//
+
+  // Collect all call graph paths from kernel `K` to callee `Callee`.
+  void collectCallGraphPaths(Function *K, Function *Callee,
+                             SmallVectorImpl<SetVector<Function *>> &CGPaths);
+
+  // TODO: What is the neat way to implement this functionality?
+  void propogateKernelId();
+
+  // Associate kernel and LDS global to unique ids and vice versa.
+  void associateKernelAndLDSGlobalToUniqueId();
+
+  // Associate kernels and LDS globals with unique integer id, and implement a
+  // mechanism to access the kernel ids within callees at run time.
+  void prepareForLowering();
+
+  //===--------------------------------------------------------------------===//
+  // Methods which aid in creating the various `map` data structures.
+  //===--------------------------------------------------------------------===//
+
+  // Associate each kernel K with LDS globals which are being accessed by K
+  // and/or by the callees of K.
+  void createKernelToLDSGlobalsMap();
+
+  // Traverse through the call graph nodes associated with the callees of
+  // current caller, and push them into stack.
+  void pushCallGraphNodes(CallGraphNode *CGNode,
+                          SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+                          SmallVectorImpl<CallBase *> &CallBaseStack);
+
+  // The call site associated with `CGNode` is a "direct call site", and is
+  // associated with a single callee, say,`Callee` represented by `CGNode`. Add
+  // `Callee` to `CalleeSet`, and push callees of `Callee` to `CGNodeStack` to
+  // further explore DFS search.
+  void collectCalleeAssociatedWithDirectCallSite(
+      CallGraphNode *CGNode, SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+      SmallVectorImpl<CallBase *> &CallBaseStack,
+      SmallPtrSetImpl<Function *> &CalleeSet);
+
+  // The call site `CB` is an "indirect call site". Resolve `CB` to a set of
+  // potential callees.
+  void associateIndirectCallSiteWithPotentialCallees(
+      CallGraphNode *CGNode, CallBase *CB,
+      SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+      SmallVectorImpl<CallBase *> &CallBaseStack,
+      SmallPtrSetImpl<Function *> &CalleeSet);
+
+  // Collect callee(s) associated with call site `CB`. If `CB` is a `direct`
+  // call site, then there is only one callee associated with it, collect it. If
+  // `CB` is an `indirect` call site, then potentially there could be more than
+  // one callee associated with it, collect all of them.
+  bool collectCallees(CallGraphNode *CGNode, CallBase *CB,
+                      SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+                      SmallVectorImpl<CallBase *> &CallBaseStack,
+                      SmallPtrSet<CallBase *, 16> &VisitedIndirectCallSites,
+                      SmallPtrSetImpl<Function *> &CalleeSet);
+
+  // Traverse `CallGraph` starting from the `CallGraphNode` associated with each
+  // kernel `K` in DFS manner and collect all the callees which are reachable
+  // from K (including indirectly called callees).
+  void createKernelToCalleesMap();
+
+  // Associate each kernel/function with the LDS globals which are being
+  // accessed within them.
+  void createAccessorToLDSGlobalsMap();
+
+  // For each `LDS`, recursively visit its user list and find all those
+  // kernels/functions within which the `LDS` is being accessed.
+  void createLDSGlobalToAccessorsMap();
+
+  // For each kernel `K`, collect LDS globals which are being accessed during
+  // the execution of `K`.
+  bool collectPerKernelAccessibleLDSGlobals();
+
+  //===--------------------------------------------------------------------===//
+  // Methods which aid in creating the various `set` data structures.
+  //===--------------------------------------------------------------------===//
+
+  // Collect all the amdgpu kernels defined within the current module.
+  bool collectKernels();
+
+  // Collect all the (static) LDS globals defined within the current module.
+  bool collectLDSGlobals();
+
+  // Collect functions whose address is taken within the module.
+  void collectAddressTakenFunctions();
+};
+
+// Convert `ConstantExpr CE` to `Instruction`, and update users of `CE` to use
+// `Instruction`.
+Instruction *LowerLDSGlobalImpl::replaceConstExprByInst(ConstantExpr *CE) {
+  // FIXME: This looks like a hack to me, but, is there any better way of
+  // handling `ConstantExprs`? I have no idea at the moment, need to revisit it
+  // later.
+  for (auto *U : CE->users()) {
+    auto *I = dyn_cast<Instruction>(U);
+
+    if (!I)
+      I = replaceConstExprByInst(dyn_cast<ConstantExpr>(U));
+
+    if (I) {
+      auto *NI = CE->getAsInstruction();
+      NI->insertBefore(I);
+      unsigned Ind = 0;
+      for (Use &UU : I->operands()) {
+        Value *V = UU.get();
+        if (V == CE)
+          I->setOperand(Ind, NI);
+        ++Ind;
+      }
+      return NI;
+    }
+  }
+
+  return nullptr;
+}
+
+// Within instruction `I` replace the use(s) of `LDS` by `LayoutAccessInst`.
+void LowerLDSGlobalImpl::replaceInstWhichUsesLDS(
+    Instruction *I, GlobalVariable *LDS, Value *LayoutAccessInst,
+    SmallPtrSetImpl<Instruction *> &ToBeErasedInsts) {
+  // Create clone of `I`, say, it is `NewI`. Within `NewI`, replace the use(s)
+  // of `LDS` by `LayoutAccessor`.
+  //
+  // FIXME: Instruction cloning is not required, fix it.
+  Instruction *NewI = I->clone();
+  unsigned Ind = 0;
+  for (Use &UU : NewI->operands()) {
+    Value *V = UU.get();
+    if (V == LDS)
+      NewI->setOperand(Ind, LayoutAccessInst);
+    ++Ind;
+  }
+
+  // Insert `NewI` just before `I`, replace all uses of `I` by `NewI` and mark
+  // `I` as `to be erased` instruction.
+  NewI->insertBefore(I);
+  NewI->copyMetadata(*I);
+  I->replaceAllUsesWith(NewI);
+  ToBeErasedInsts.insert(I);
+}
+
+// Replace all the users of original LDS global `LDS` with new ones which access
+// corresponding LDS layouts.
+void LowerLDSGlobalImpl::replaceAllUsersOfLDS(
+    GlobalVariable *LDS,
+    ValueMap<Function *, Instruction *> &FunctionToLDSLayoutAccessInst) {
+  // Keep track of all the erased to be instructions.
+  SmallPtrSet<Instruction *, 16> ToBeErasedInsts;
+
+  // Traverse through each use `U` of `LDS`, create a new one which replaces `U`
+  // and accordingly replace `U`.
+  for (auto *U : LDS->users()) {
+    // `U` may be using `LDS`, but 'U` itself is not used anywhere, ignore `U`.
+    if (!U->getNumUses())
+      continue;
+
+    if (auto *I = dyn_cast<Instruction>(U)) {
+      // User is an Instruction, accordingly handle it.
+      auto *LayoutAccessInst =
+          FunctionToLDSLayoutAccessInst[I->getParent()->getParent()];
+      replaceInstWhichUsesLDS(I, LDS, LayoutAccessInst, ToBeErasedInsts);
+    } else if (auto *CE = dyn_cast<ConstantExpr>(U)) {
+      // User is a ConstantExpr, accordingly handle it.
+      auto *I = replaceConstExprByInst(CE);
+      auto *LayoutAccessInst =
+          FunctionToLDSLayoutAccessInst[I->getParent()->getParent()];
+      replaceInstWhichUsesLDS(I, LDS, LayoutAccessInst, ToBeErasedInsts);
+      CE->removeDeadConstantUsers();
+    } else
+      llvm_unreachable("Not Implemented."); // TODO: What else is missing?
+  }
+
+  // Erase all the instructions which are got replaced by new ones.
+  for (auto *I : ToBeErasedInsts)
+    I->eraseFromParent();
+}
+
+// Add instructions to access LDS layouts corresponding to LDS global `LDS`.
+bool LowerLDSGlobalImpl::accessLDSLayouts(
+    GlobalVariable *LDS,
+    ValueMap<Function *, Instruction *> &FunctionToLDSLayoutAccessInst) {
+  // LDS global is not used anywhere? ignore it.
+  if (!contains(LDSGlobalToAccessors, LDS))
+    return false;
+
+  for (auto *F : LDSGlobalToAccessors[LDS]) {
+    // Get instruction insertion point.
+    Instruction *EI = AMDGPU::isModuleEntryFunctionCC(F->getCallingConv())
+                          ? KernelToInstInsertPt[F]
+                          : &(*F->getEntryBlock().getFirstInsertionPt());
+
+    // Get indices to access LDS offset table.
+    // FIXME: explicit cast
+    Value *KI = AMDGPU::isModuleEntryFunctionCC(F->getCallingConv())
+                    ? (Value *)Constant::getIntegerValue(
+                          Type::getInt32Ty(Ctx), APInt(32, KernelToID[F]))
+                    : (Value *)F->getArg(F->arg_size() - 1);
+    Value *LI = Constant::getIntegerValue(Type::getInt32Ty(Ctx),
+                                          APInt(32, LDSToID[LDS]));
+
+    // Insert GEP instruction to access the LDS offset table at the address
+    // ((LDSOffsetTable + KI) + LI) which is of type `i32*`.
+    Value *Indices1[] = {Constant::getNullValue(Type::getInt32Ty(Ctx)), KI, LI};
+    Instruction *GEPI1 = GetElementPtrInst::CreateInBounds(
+        LDSOffsetTable->getValueType(), LDSOffsetTable, Indices1,
+        PrefixStr + Twine("gep."), EI);
+
+    // Insert LOAD instruction to load the `offset` value from LDS offset table
+    // at the address ((LDSOffsetTable + KI) + LI) which is of type `i32`.
+    // FIXME: getPointerElementType() sould not be used.
+    Instruction *LoadI = new LoadInst(GEPI1->getType()->getPointerElementType(),
+                                      GEPI1, PrefixStr + Twine("load."), EI);
+
+    // Insert GEP instruction to access the starting address of LDS layout
+    // corresponding to kernel represented by the kernel id `KI` which is of
+    // type `i8*`.
+    Value *Indices2[] = {Constant::getNullValue(Type::getInt32Ty(Ctx)), KI};
+    Instruction *GEPI2 = GetElementPtrInst::CreateInBounds(
+        LDSLayouts->getValueType(), LDSLayouts, Indices2,
+        PrefixStr + Twine("gep."), EI);
+
+    // Insert GEP instruction to access the base address corresponding to LDS
+    // global `LDS` within above LDS layout which is of type 'i8*`.
+    // FIXME: getPointerElementType() sould not be used.
+    Instruction *GEPI3 = GetElementPtrInst::CreateInBounds(
+        GEPI2->getType()->getPointerElementType(), GEPI2, LoadI,
+        PrefixStr + Twine("gep."), EI);
+
+    // Insert BITCAST instruction to cast `above base address` from 'i8*` to
+    // `ldstype*` where `ldstype` is the type of original LDS global `LDS`.
+    Instruction *CastI =
+        new BitCastInst(GEPI3, LDS->getType(), PrefixStr + Twine("cast."), EI);
+
+    // Save the above instruction `CastI` which replaces all uses of `LDS`
+    // within `F`.
+    FunctionToLDSLayoutAccessInst[F] = CastI;
+  }
+
+  return true;
+}
+
+// Add instructions within kernels to initialize LDS layouts array to hold the
+// starting address of LDS layouts.
+void LowerLDSGlobalImpl::initializeArrayOfLDSLayouts() {
+  for (auto *K : Kernels) {
+    // Kernel `K` does not associate with LDS layout? skip it.
+    if (!contains(KernelToLDSLayout, K))
+      continue;
+
+    // Insert instructions at the beginning of the entry basic block of the
+    // kernel `K` which initialize the global array `LDSLayouts` to hold the
+    // starting address of the LDS layout associated with `K` at LDSLayouts[ID]
+    // where ID is kernel id of `K`.
+    Instruction *EI = &(*(K->getEntryBlock().getFirstInsertionPt()));
+
+    // Insert GEP instruction which access the address `LDSLayout + 0`, say, the
+    // result is `GEP1` which is of type `i8*`.
+    GlobalVariable *LDSLayout = KernelToLDSLayout[K];
+    Value *Indices1[] = {
+        Constant::getNullValue(Type::getInt32Ty(Ctx)),
+        Constant::getIntegerValue(Type::getInt32Ty(Ctx), APInt(32, 0))};
+    Instruction *GEP1 = GetElementPtrInst::CreateInBounds(
+        LDSLayout->getValueType(), const_cast<GlobalVariable *>(LDSLayout),
+        Indices1, PrefixStr + Twine("gep."), EI);
+
+    // Insert GEP instruction which access the address `LDSLayouts + ID`, say,
+    // the result is `GEP2` which is of type `i8**`.
+    Value *Indices2[] = {Constant::getNullValue(Type::getInt32Ty(Ctx)),
+                         Constant::getIntegerValue(Type::getInt32Ty(Ctx),
+                                                   APInt(32, KernelToID[K]))};
+    Instruction *GEP2 = GetElementPtrInst::CreateInBounds(
+        LDSLayouts->getValueType(), const_cast<GlobalVariable *>(LDSLayouts),
+        Indices2, PrefixStr + Twine("gep."), EI);
+
+    // Insert STORE instruction which stores `GEP1` at `GEP2`.
+    new StoreInst(GEP1, GEP2, EI);
+
+    // Save the instruction insertion point which will be later required when
+    // it is necessary to insert LDS layout access instructions within kernel.
+    KernelToInstInsertPt[K] = EI;
+  }
+}
+
+// Insert various required instructions within kernels and within callees which
+// aid in accessing of newly inserted LDS layouts.
+void LowerLDSGlobalImpl::insertNewInstructions() {
+  // Add instructions within kernels to initialize LDS layouts array to hold the
+  // starting address of LDS layouts.
+  initializeArrayOfLDSLayouts();
+
+  // Add instructions to access LDS layouts and accordingly replace original
+  // instructions which use original LDS globals.
+  for (auto *LDS : LDSGlobals) {
+    ValueMap<Function *, Instruction *> FunctionToLDSLayoutAccessInst;
+
+    // LDS global is not used anywhere? ignore it.
+    if (!accessLDSLayouts(LDS, FunctionToLDSLayoutAccessInst))
+      continue;
+
+    // Replace all the users of LDS global `LDS` with new ones which access
+    // corresponding LDS layouts.
+    replaceAllUsersOfLDS(LDS, FunctionToLDSLayoutAccessInst);
+  }
+}
+
+// Create 2D LDS offset table which will be referenced at run time to access the
+// LDS specific offset within kernel specific LDS layout.
+void LowerLDSGlobalImpl::constructLDSOffsetTable() {
+  // Get type of 2D LDS offset table.
+  auto *EleTy = Type::getInt32Ty(Ctx);
+  auto *Arr1DTy = ArrayType::get(EleTy, LDSGlobals.size());
+  auto *Arr2DTy = ArrayType::get(Arr1DTy, Kernels.size());
+
+  // Create offset initialization list.
+  SmallVector<llvm::Constant *, 16> Init2DValues;
+
+  for (unsigned K = 0; K < Kernels.size(); ++K) {
+    SmallVector<llvm::Constant *, 16> Init1DValues;
+    auto *Kernel = IDToKernel[K];
+
+    for (unsigned L = 0; L < LDSGlobals.size(); ++L) {
+      auto *LDS = IDToLDS[L];
+      auto Offset = contains(KernelToLDSToOffset, Kernel)
+                        ? contains(KernelToLDSToOffset[Kernel], LDS)
+                              ? KernelToLDSToOffset[Kernel][LDS]
+                              : -1
+                        : -1;
+      Init1DValues.push_back(
+          Constant::getIntegerValue(EleTy, APInt(32, Offset)));
+    }
+
+    Init2DValues.push_back(ConstantArray::get(Arr1DTy, Init1DValues));
+  }
+
+  // Create 2D LDS offset table which is initialized with offsets.
+  LDSOffsetTable = new GlobalVariable(
+      M, Arr2DTy, false, GlobalValue::InternalLinkage,
+      ConstantArray::get(Arr2DTy, Init2DValues),
+      Twine("__llvm.amdgcn.lds.offset.table__"), nullptr,
+      GlobalVariable::NotThreadLocal, AMDGPUAS::CONSTANT_ADDRESS);
+
+  // Set proper alignment.
+  LDSOffsetTable->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+  LDSOffsetTable->setAlignment(getAlign(LDSOffsetTable));
+}
+
+// Create 1D array which holds LDS layouts which will be referenced at run time
+// to access kernel specific LDS layout.
+void LowerLDSGlobalImpl::constructArrayOfLDSLayouts() {
+  auto *ArrTy = ArrayType::get(
+      PointerType::get(IntegerType::get(Ctx, 8), AMDGPUAS::LOCAL_ADDRESS),
+      Kernels.size());
+
+  LDSLayouts = new GlobalVariable(
+      M, ArrTy, false, GlobalValue::InternalLinkage, UndefValue::get(ArrTy),
+      Twine("__llvm.amdgcn.unified.lds.layouts__"), nullptr,
+      GlobalVariable::NotThreadLocal, AMDGPUAS::CONSTANT_ADDRESS);
+
+  LDSLayouts->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+  LDSLayouts->setAlignment(getAlign(LDSLayouts));
+}
+
+// Create a new LDS layout corresponding to kernel `K` whose size will be of
+// `LayoutSize`.
+GlobalVariable *LowerLDSGlobalImpl::constructLDSLayout(Function *K,
+                                                       uint32_t LayoutSize) {
+  auto *LayoutType = ArrayType::get(IntegerType::get(Ctx, 8), LayoutSize);
+
+  auto *LDSLayout = new GlobalVariable(
+      M, LayoutType, false, GlobalValue::InternalLinkage,
+      UndefValue::get(LayoutType),
+      Twine("__llvm.amdgcn.unified.lds.layout.") + Twine(KernelToID[K]) +
+          Twine("__"),
+      nullptr, GlobalVariable::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS);
+
+  LDSLayout->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+  LDSLayout->setAlignment(getAlign(LDSLayout));
+
+  return LDSLayout;
+}
+
+// Given a set of LDS globals, associate each LDS from the set to a
+// corresponding offset within the kernel specific LDS layout.
+uint32_t LowerLDSGlobalImpl::computeOffsets(
+    SmallPtrSetImpl<GlobalVariable *> &LDSSet,
+    std::map<GlobalVariable *, uint32_t> &LDSToOffSet) {
+  // Sort `LDSSet` by alignment in descending order to minimize the padding
+  // required. On ties, sort by type allocation size, again in descending order,
+  // and finally, by name in lexicographical order (in ascending order) when
+  // type allocation sizes are same.
+  std::vector<GlobalVariable *> SortedLDSSet(LDSSet.begin(), LDSSet.end());
+  llvm::stable_sort(SortedLDSSet,
+                    [&](GlobalVariable *LHS, GlobalVariable *RHS) -> bool {
+                      auto ALHS = getAlign(LHS);
+                      auto ARHS = getAlign(RHS);
+                      if (ALHS != ARHS)
+                        return ALHS > ARHS;
+
+                      auto SLHS = DL.getTypeAllocSize(LHS->getValueType());
+                      auto SRHS = DL.getTypeAllocSize(RHS->getValueType());
+                      if (SLHS != SRHS)
+                        return SLHS > SRHS;
+
+                      return LHS->getName() < RHS->getName();
+                    });
+
+  // Compute `offset` for each LDS global from `SortedLDSSet`.
+  // FIXME: Alignment is not yet handled properly.
+  uint32_t CurOffset = 0;
+  for (auto *LDS : SortedLDSSet) {
+    LDSToOffSet[LDS] = CurOffset;
+    CurOffset += DL.getTypeAllocSize(LDS->getValueType()).getFixedValue();
+  }
+
+  return CurOffset;
+}
+
+// Create unified LDS layouts for each kernel.
+void LowerLDSGlobalImpl::constructLDSLayouts() {
+  // Create a new LDS layout corresponding to each kernel `K`, and associate
+  // each kernel to a map which maps LDS globals to corresponding offsets.
+  // FIXME: Alignment is not yet handled properly.
+  for (auto *K : Kernels) {
+    // Kernel `K` does not associate with any LDS global? skip it.
+    if (!contains(KernelToLDSGlobals, K))
+      continue;
+
+    // For the set of LDS globals associated with kernel `K`, map each LDS
+    // from the set to a corresponding offset within the kernel specific LDS
+    // layout, and also get the grand total size of LDS layout.
+    std::map<GlobalVariable *, uint32_t> LDSToOffSet;
+    auto LayoutSize = computeOffsets(KernelToLDSGlobals[K], LDSToOffSet);
+    KernelToLDSToOffset[K] = LDSToOffSet;
+
+    // Create a new LDS layout corresponding to kernel `K` whose size will be of
+    // `LayoutSize`.
+    KernelToLDSLayout[K] = constructLDSLayout(K, LayoutSize);
+  }
+}
+
+// Insert various required globals (including per kernel specific unified LDS
+// layouts) into the module.
+void LowerLDSGlobalImpl::insertNewGlobals() {
+  // Create unified LDS layouts for each kernel.
+  constructLDSLayouts();
+
+  // Create 1D array which holds LDS layouts which will be referenced at run
+  // time to access kernel specific LDS layout.
+  constructArrayOfLDSLayouts();
+
+  // Create 2D LDS offset table which will be referenced at run time to access
+  // the LDS specific offset within kernel specific LDS layout.
+  constructLDSOffsetTable();
+}
+
+// TODO: What is the neat way to implement this functionality?
+void LowerLDSGlobalImpl::propogateKernelId() {
+  // TODO
+}
+
+// Associate kernel and LDS global to unique ids and vice versa.
+void LowerLDSGlobalImpl::associateKernelAndLDSGlobalToUniqueId() {
+  // Associate each kernel to unique integer and vice versa.
+  uint32_t KNum = 0;
+  for (auto *K : Kernels) {
+    KernelToID[K] = KNum;
+    IDToKernel[KNum] = K;
+    ++KNum;
+  }
+
+  // Associate each LDS to unique integer and vice versa.
+  uint32_t LNum = 0;
+  for (auto *LDSGlobal : LDSGlobals) {
+    LDSToID[LDSGlobal] = LNum;
+    IDToLDS[LNum] = LDSGlobal;
+    ++LNum;
+  }
+}
+
+// Associate kernels and LDS globals with unique integer id, and implement a
+// mechanism to access the kernel ids within callees at run time.
+void LowerLDSGlobalImpl::prepareForLowering() {
+  // Associate kernel and LDS global to unique ids and vice versa.
+  associateKernelAndLDSGlobalToUniqueId();
+
+  // TODO: What is the neat way to implement this functionality?
+  propogateKernelId();
+}
+
+// Associate each kernel K with LDS globals which are being accessed by K and/or
+// by the callees of K.
+void LowerLDSGlobalImpl::createKernelToLDSGlobalsMap() {
+  for (auto *K : Kernels) {
+    SmallPtrSet<GlobalVariable *, 8> LDSSet;
+
+    // Collect all those LDS globals which are being accessed by kernel K
+    // itself.
+    if (contains(AccessorToLDSGlobals, K))
+      LDSSet.insert(AccessorToLDSGlobals[K].begin(),
+                    AccessorToLDSGlobals[K].end());
+
+    // Collect all those LDS globals which are being accessed by the callees of
+    // kernel K.
+    for (auto *Callee : KernelToCallees[K]) {
+      if (contains(AccessorToLDSGlobals, Callee))
+        LDSSet.insert(AccessorToLDSGlobals[Callee].begin(),
+                      AccessorToLDSGlobals[Callee].end());
+    }
+
+    if (!LDSSet.empty())
+      KernelToLDSGlobals[K] = LDSSet;
+  }
+}
+
+// Traverse through the call graph nodes associated with the callees of current
+// caller, and push them into stack.
+void LowerLDSGlobalImpl::pushCallGraphNodes(
+    CallGraphNode *CGNode, SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+    SmallVectorImpl<CallBase *> &CallBaseStack) {
+  for (auto GI = CGNode->begin(), GE = CGNode->end(); GI != GE; ++GI) {
+    auto *CGN = GI->second;
+    auto *CB = cast<CallBase>(GI->first.getValue());
+    CGNodeStack.push_back(CGN);
+    CallBaseStack.push_back(CB);
+  }
+}
+
+// The call site associated with `CGNode` is a "direct call site", and is
+// associated with a single callee, say,`Callee` represented by `CGNode`. Add
+// `Callee` to `CalleeSet`, and push callees of `Callee` to `CGNodeStack` to
+// further explore DFS search.
+void LowerLDSGlobalImpl::collectCalleeAssociatedWithDirectCallSite(
+    CallGraphNode *CGNode, SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+    SmallVectorImpl<CallBase *> &CallBaseStack,
+    SmallPtrSetImpl<Function *> &CalleeSet) {
+  auto *Callee = CGNode->getFunction();
+  assert(Callee && "Expected a valid callee associated with call site");
+  assert(!Callee->isDeclaration() &&
+         "Expected callee defintion since this pass assumes that caller and "
+         "callee appear in same translation unit");
+  CalleeSet.insert(Callee);
+  pushCallGraphNodes(CGNode, CGNodeStack, CallBaseStack);
+}
+
+// The call site `CB` is an "indirect call site". Resolve `CB` to a set of
+// potential callees.
+void LowerLDSGlobalImpl::associateIndirectCallSiteWithPotentialCallees(
+    CallGraphNode *CGNode, CallBase *CB,
+    SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+    SmallVectorImpl<CallBase *> &CallBaseStack,
+    SmallPtrSetImpl<Function *> &CalleeSet) {
+  if (auto *MD = CB->getMetadata(LLVMContext::MD_callees)) {
+    // The metadata "!callee" is available at the indirect call site `CB`, which
+    // means, all the potential target callees for the call site `CB` is
+    // successfully resolved at compile time. So, push them into stack so that
+    // they will be handled just like direct callees when they are eventually
+    // poped out.
+    for (const auto &Op : MD->operands()) {
+      CGNodeStack.push_back(CG[mdconst::extract_or_null<Function>(Op)]);
+      CallBaseStack.push_back(CB);
+    }
+  } else {
+    // The metadata "!callee" is *NOT* available at the indirect call site `CB`,
+    // which means, `CB` has *NO* information about potential target callees.
+    // The simplest possible *SAFE* assumption that we can make here is to
+    // consider all those "address taken" functions whose singature matches with
+    // that of the call site `CB`, and assume that all these signature matched
+    // "address taken" functions are possible potential callees. Thus, push all
+    // the signature matched "address taken" functions into stack so that they
+    // will be handled just like direct callees when they are eventually poped
+    // out.
+    auto *CBFTy = CB->getFunctionType();
+    for (auto *CGN : AddressTakenSet) {
+      auto *F = CGN->getFunction();
+      assert(F && "Expected a valid address taken function");
+      auto *ADFTy = F->getFunctionType();
+      if (ADFTy == CBFTy) {
+        CGNodeStack.push_back(CGN);
+        CallBaseStack.push_back(CB);
+      }
+    }
+  }
+}
+
+// Collect callee(s) associated with call site `CB`. If `CB` is a `direct` call
+// site, then there is only one callee associated with it, collect it. If `CB`
+// is an `indirect` call site, then potentially there could be more than one
+// callee associated with it, collect all of them.
+bool LowerLDSGlobalImpl::collectCallees(
+    CallGraphNode *CGNode, CallBase *CB,
+    SmallVectorImpl<CallGraphNode *> &CGNodeStack,
+    SmallVectorImpl<CallBase *> &CallBaseStack,
+    SmallPtrSet<CallBase *, 16> &VisitedIndirectCallSites,
+    SmallPtrSetImpl<Function *> &CalleeSet) {
+  // Return `true` if the collees are actually collected.
+  bool CollectedCollees = true;
+
+  if (!CB->getCalledFunction()) {
+    // Call site `CB` is an indirect call site. But, if the `CGNode` has a
+    // function defintion, say, `Callee`, associated with it, which means, we
+    // have already had visited `CB` earlier, and we had resovled it to a set of
+    // pontential callees, and `Callee` is one among them. Collect `Callee` just
+    // the way direct callee is collected. Otherwise, `CB` is encoutered for
+    // first time. Resolve it to a set of potential callees before collecting
+    // them.
+    if (!VisitedIndirectCallSites.contains(CB)) {
+      // Indirect call site `CB` is encoutered first time, resolve it to a set
+      // of potential callees.
+      associateIndirectCallSiteWithPotentialCallees(CGNode, CB, CGNodeStack,
+                                                    CallBaseStack, CalleeSet);
+
+      // Mark that `CB` is visited.
+      VisitedIndirectCallSites.insert(CB);
+
+      // We are only resolved indirect call site `CB` to a set of potential
+      // callees here, and they are not collected yet.
+      CollectedCollees = false;
+    } else {
+      // Indirect call site `CB` is already been resolved to a set of potential
+      // callees during its first visit. The callee represented by `CGNode` is
+      // one among them. Collect it just the way direct callee is collected.
+      collectCalleeAssociatedWithDirectCallSite(CGNode, CGNodeStack,
+                                                CallBaseStack, CalleeSet);
+    }
+  } else {
+    // Call site `CB` is a direct call site. Collect a single callee which is
+    // associated with it.
+    collectCalleeAssociatedWithDirectCallSite(CGNode, CGNodeStack,
+                                              CallBaseStack, CalleeSet);
+  }
+
+  return CollectedCollees;
+}
+
+// Traverse `CallGraph` starting from the `CallGraphNode` associated with each
+// kernel `K` in DFS manner and collect all the callees which are reachable from
+// K (including indirectly called callees).
+//
+// FIXME: Can we "really" get rid of this function, and is it that llvm
+// "CallGraph" infra structure already provides this functionality, especially,
+// by including indirect callees within SCC?
+void LowerLDSGlobalImpl::createKernelToCalleesMap() {
+  for (auto *K : Kernels) {
+    auto *KernCGNode = CG[K];
+    SmallVector<CallGraphNode *, 16> CGNodeStack;
+    SmallVector<CallBase *, 16> CallBaseStack;
+    SmallPtrSet<CallGraphNode *, 16> VisitedCGNodes;
+    SmallPtrSet<CallBase *, 16> VisitedIndirectCallSites;
+    SmallPtrSet<Function *, 8> CalleeSet;
+
+    // Push the `CallGraphNode` associated with all the callees of the kernel`K`
+    // into into `CGNodeStack`, and the corresponding call sites into
+    // `CallBaseStack`.
+    pushCallGraphNodes(KernCGNode, CGNodeStack, CallBaseStack);
+
+    // Continue DFS search until no more call graph nodes to handle.
+    while (!CGNodeStack.empty()) {
+      assert(CGNodeStack.size() == CallBaseStack.size() &&
+             "Stack holding CallBase pointers is currupted");
+      auto *CGNode = CGNodeStack.pop_back_val();
+      auto *CB = CallBaseStack.pop_back_val();
+
+      // `CGNode` is already visited and handled, ignore it and proceed to next
+      // one.
+      if (!VisitedCGNodes.insert(CGNode).second)
+        continue;
+
+      // Collect callee(s) associated with call site `CB`. If `CB` is a `direct`
+      // call site, then there is only one callee associated with it, collect
+      // it. If `CB` is an `indirect` call site, then potentially there could be
+      // more than one callee associated with it, collect all of them.
+      auto CallSiteHandledIsIndirect =
+          collectCallees(CGNode, CB, CGNodeStack, CallBaseStack,
+                         VisitedIndirectCallSites, CalleeSet);
+
+      // The call site `CB` which is being handled just now is an indirect call
+      // site. Since the indirect call site does not bind to any particular set
+      // of callees, and since the `CGNode` is one and the same for all the
+      // function pointers which have same signature, we should *NOT* assume
+      // that `CGNode` is visited unlike in case of direct call site.
+      if (CallSiteHandledIsIndirect)
+        VisitedCGNodes.erase(CGNode);
+    }
+
+    assert(CallBaseStack.empty() &&
+           "Stack holding CallBase pointers is currupted");
+
+    KernelToCallees[K] = CalleeSet;
+  }
+}
+
+// Associate each kernel/function with the LDS globals which are being accessed
+// within them.
+void LowerLDSGlobalImpl::createAccessorToLDSGlobalsMap() {
+  for (auto LI = LDSGlobalToAccessors.begin(), LE = LDSGlobalToAccessors.end();
+       LI != LE; ++LI) {
+    for (auto *A : LI->second) {
+      if (!contains(AccessorToLDSGlobals, A)) {
+        SmallPtrSet<GlobalVariable *, 8> LDSSet;
+        LDSSet.insert(LI->first);
+        AccessorToLDSGlobals[A] = LDSSet;
+      } else
+        AccessorToLDSGlobals[A].insert(LI->first);
+    }
+  }
+}
+
+// For each `LDS`, recursively visit its user list and find all those
+// kernels/functions within which the `LDS` is being accessed.
+void LowerLDSGlobalImpl::createLDSGlobalToAccessorsMap() {
+  for (auto *LDS : LDSGlobals) {
+    assert(!LDS->user_empty() &&
+           "LDS user list cannot be empty since it must have been successfully "
+           "defined within either kernel or function");
+
+    SmallPtrSet<Function *, 8> LDSAccessors;
+    SmallVector<User *, 16> UserStack(LDS->users());
+    SmallPtrSet<User *, 16> VisitedUsers;
+
+    while (!UserStack.empty()) {
+      auto *U = UserStack.pop_back_val();
+      if (!VisitedUsers.insert(U).second)
+        continue;
+
+      auto *I = dyn_cast<Instruction>(U);
+
+      // If `U` is not an `Instruction`, then it should be a `Constant` which is
+      // nested within an `Instruction`. Push-back users of `U`, and continue
+      // further exploring the stack until an `Instruction` is found.
+      if (!I) {
+        assert(isa<Constant>(U) && "Expected a constant expression");
+        append_range(UserStack, U->users());
+        continue;
+      }
+
+      // We have successfully found a kernel/function within which the `LDS` is
+      // being accessed, insert it into `LDSAccessors` set.
+      LDSAccessors.insert(I->getParent()->getParent());
+    }
+
+    LDSGlobalToAccessors[LDS] = LDSAccessors;
+  }
+}
+
+// For each kernel `K`, collect LDS globals which are being accessed during the
+// execution of `K`.
+bool LowerLDSGlobalImpl::collectPerKernelAccessibleLDSGlobals() {
+  // Associate each LDS with the kernels/functions within which the LDS is being
+  // accessed.
+  createLDSGlobalToAccessorsMap();
+
+  // Associate each kernel/function with the LDS globals which are being
+  // accessed within them.
+  createAccessorToLDSGlobalsMap();
+
+  // Associate each kernel K with callees which are reachable from K (including
+  // indirectly called callees).
+  createKernelToCalleesMap();
+
+  // Associate each kernel K with LDS globals which are being accessed by K
+  // and/or by the callees of K.
+  createKernelToLDSGlobalsMap();
+
+  // If *none* of the kernels associate with any LDS globals, then nothing do.
+  return !KernelToLDSGlobals.empty();
+}
+
+// Collect all the amdgpu kernels defined within the current module.
+bool LowerLDSGlobalImpl::collectKernels() {
+  for (auto &F : M.functions()) {
+    if (!F.isDeclaration() &&
+        AMDGPU::isModuleEntryFunctionCC(F.getCallingConv()))
+      Kernels.insert(&F);
+  }
+
+  return !Kernels.empty();
+}
+
+// Collect all the (static) LDS globals defined within the current module.
+bool LowerLDSGlobalImpl::collectLDSGlobals() {
+  for (auto &GV : M.globals()) {
+    if (GV.getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS &&
+        !GV.isDeclaration() && !GV.getType()->isEmptyTy())
+      LDSGlobals.insert(&GV);
+  }
+
+  return !LDSGlobals.empty();
+}
+
+// Collect functions whose address is taken within the module.
+void LowerLDSGlobalImpl::collectAddressTakenFunctions() {
+  auto *ExternalCallingNode = CG.getExternalCallingNode();
+  for (auto GI = ExternalCallingNode->begin(), GE = ExternalCallingNode->end();
+       GI != GE; ++GI) {
+    auto *CGN = GI->second;
+    auto *F = CGN->getFunction();
+    // FIXME: Anything else need to be excluded?
+    if (!F || F->isDeclaration() ||
+        AMDGPU::isModuleEntryFunctionCC(F->getCallingConv()))
+      continue;
+    AddressTakenSet.insert(CGN);
+  }
+}
+
+// Entry-point function.
+bool LowerLDSGlobalImpl::lower() {
+  // If there are *no* LDS globals defined within the module, or if there are
+  // *no* kernels defined within the module, or if there exist *no* kernel
+  // *execution* which accesses LDS globals at run time, then nothing to do.
+  if (!collectLDSGlobals() || !collectKernels() ||
+      !collectPerKernelAccessibleLDSGlobals())
+    return false;
+
+  // Associate kernels and LDS globals with unique integer id, and implement a
+  // mechanism to access the kernel ids within callees at run time.
+  prepareForLowering();
+
+  // Insert various required globals (including per kernel specific unified LDS
+  // layouts) into the module.
+  insertNewGlobals();
+
+  // Insert various required instructions within kernels and within callees
+  // which aid accessing of newly inserted LDS layouts.
+  insertNewInstructions();
+
+  return true;
+}
+
+class AMDGPULowerLDSGlobal : public ModulePass {
+public:
+  static char ID;
+
+  AMDGPULowerLDSGlobal() : ModulePass(ID) {
+    initializeAMDGPULowerLDSGlobalPass(*PassRegistry::getPassRegistry());
+  }
+
+  bool runOnModule(Module &M) override;
+};
+
+} // namespace
+
+char AMDGPULowerLDSGlobal::ID = 0;
+char &llvm::AMDGPULowerLDSGlobalID = AMDGPULowerLDSGlobal::ID;
+
+INITIALIZE_PASS(AMDGPULowerLDSGlobal, "amdgpu-lower-lds-global",
+                "Lower LDS Global Variables", false /*only look at the cfg*/,
+                false /*analysis pass*/)
+
+bool AMDGPULowerLDSGlobal::runOnModule(Module &M) {
+  LowerLDSGlobalImpl LDSLowerer{M};
+  return LDSLowerer.lower();
+}
+
+ModulePass *llvm::createAMDGPULowerLDSGlobalPass() {
+  return new AMDGPULowerLDSGlobal();
+}
+
+PreservedAnalyses AMDGPULowerLDSGlobalPass::run(Module &M,
+                                                ModuleAnalysisManager &AM) {
+  LowerLDSGlobalImpl LDSLowerer{M};
+  LDSLowerer.lower();
+  return PreservedAnalyses::all();
+}

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h

protected:
 
   StringRef getGPUName(const Function &F) const;
   StringRef getFeatureString(const Function &F) const;
 
 public:
   static bool EnableLateStructurizeCFG;
   static bool EnableFunctionCalls;
   static bool EnableFixedFunctionABI;
+  static bool EnableLDSGlobalLowering;
 
   AMDGPUTargetMachine(const Target &T, const Triple &TT, StringRef CPU,
                       StringRef FS, TargetOptions Options,
                       Optional<Reloc::Model> RM, Optional<CodeModel::Model> CM,
                       CodeGenOpt::Level OL);
   ~AMDGPUTargetMachine() override;
 
   const TargetSubtargetInfo *getSubtargetImpl() const;

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

static cl::opt<bool> EnableScalarIRPasses(
   cl::desc("Enable scalar IR passes"),
   cl::init(true),
   cl::Hidden);
 
 static cl::opt<bool> EnableStructurizerWorkarounds(
     "amdgpu-enable-structurizer-workarounds",
     cl::desc("Enable workarounds for the StructurizeCFG pass"), cl::init(true),
     cl::Hidden);
 
+static cl::opt<bool, true> EnableLDSGlobalLowering(
+    "amdgpu-enable-lds-global-lowering",
+    cl::desc("Enable LDS global variable lowering pass"),
+    cl::location(AMDGPUTargetMachine::EnableLDSGlobalLowering), cl::init(true),
+    cl::Hidden);
+
 extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
   // Register the target
   RegisterTargetMachine<R600TargetMachine> X(getTheAMDGPUTarget());
   RegisterTargetMachine<GCNTargetMachine> Y(getTheGCNTarget());
 
   PassRegistry *PR = PassRegistry::getPassRegistry();
   initializeR600ClauseMergePassPass(*PR);
   initializeR600ControlFlowFinalizerPass(*PR);
   initializeR600PacketizerPass(*PR);
extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
   initializeAMDGPUAAWrapperPassPass(*PR);
   initializeAMDGPUExternalAAWrapperPass(*PR);
   initializeAMDGPUUseNativeCallsPass(*PR);
   initializeAMDGPUSimplifyLibCallsPass(*PR);
   initializeAMDGPUPrintfRuntimeBindingPass(*PR);
   initializeGCNRegBankReassignPass(*PR);
   initializeGCNNSAReassignPass(*PR);
   initializeSIAddIMGInitPass(*PR);
+  initializeAMDGPULowerLDSGlobalPass(*PR);
 }
 
 static std::unique_ptr<TargetLoweringObjectFile> createTLOF(const Triple &TT) {
   return std::make_unique<AMDGPUTargetObjectFile>();
 }
 
 static ScheduleDAGInstrs *createR600MachineScheduler(MachineSchedContext *C) {
   return new ScheduleDAGMILive(C, std::make_unique<R600SchedStrategy>());
if (TT.getArch() == Triple::amdgcn) {
     else if (getMCSubtargetInfo()->checkFeatures("+wavefrontsize32"))
       MRI.reset(llvm::createGCNMCRegisterInfo(AMDGPUDwarfFlavour::Wave32));
   }
 }
 
 bool AMDGPUTargetMachine::EnableLateStructurizeCFG = false;
 bool AMDGPUTargetMachine::EnableFunctionCalls = false;
 bool AMDGPUTargetMachine::EnableFixedFunctionABI = false;
+bool AMDGPUTargetMachine::EnableLDSGlobalLowering = false;
 
 AMDGPUTargetMachine::~AMDGPUTargetMachine() = default;
 
 StringRef AMDGPUTargetMachine::getGPUName(const Function &F) const {
   Attribute GPUAttr = F.getFnAttribute("target-cpu");
   return GPUAttr.isValid() ? GPUAttr.getValueAsString() : getTargetCPU();
 }
 
PB.registerPipelineParsingCallback(
         if (PassName == "amdgpu-printf-runtime-binding") {
           PM.addPass(AMDGPUPrintfRuntimeBindingPass());
           return true;
         }
         if (PassName == "amdgpu-always-inline") {
           PM.addPass(AMDGPUAlwaysInlinePass());
           return true;
         }
+        if (PassName == "amdgpu-lower-lds-global") {
+          PM.addPass(AMDGPULowerLDSGlobalPass());
+          return true;
+        }
         return false;
       });
   PB.registerPipelineParsingCallback(
       [this](StringRef PassName, FunctionPassManager &PM,
              ArrayRef<PassBuilder::PipelineElement>) {
         if (PassName == "amdgpu-simplifylib") {
           PM.addPass(AMDGPUSimplifyLibCallsPass(*this));
           return true;

 void AMDGPUPassConfig::addIRPasses() {
   const AMDGPUTargetMachine &TM = getAMDGPUTargetMachine();
 
   // There is no reason to run these.
   disablePass(&StackMapLivenessID);
   disablePass(&FuncletLayoutID);
   disablePass(&PatchableFunctionID);
 
+  // We expect to run this pass as a first AMDGPU IR pass so that new
+  // instructions being added in this pass can possibly undergo further
+  // transformations via subsequent passes.
+  if (EnableLDSGlobalLowering)
+    addPass(createAMDGPULowerLDSGlobalPass());
+
   addPass(createAMDGPUPrintfRuntimeBinding());
 
   // This must occur before inlining, as the inliner will not look through
   // bitcast calls.
   addPass(createAMDGPUFixFunctionBitcastsPass());
 
   // A call to propagate attributes pass in the backend in case opt was not run.
   addPass(createAMDGPUPropagateAttributesEarlyPass(&TM));

llvm/lib/Target/AMDGPU/CMakeLists.txt

add_llvm_target(AMDGPUCodeGen
   AMDGPUGlobalISelUtils.cpp
   AMDGPULateCodeGenPrepare.cpp
   AMDGPULegalizerInfo.cpp
   AMDGPULibCalls.cpp
   AMDGPULibFunc.cpp
   AMDGPULowerIntrinsics.cpp
   AMDGPULowerKernelArguments.cpp
   AMDGPULowerKernelAttributes.cpp
+  AMDGPULowerLDSGlobal.cpp
   AMDGPUMachineCFGStructurizer.cpp
   AMDGPUMachineFunction.cpp
   AMDGPUMachineModuleInfo.cpp
   AMDGPUMacroFusion.cpp
   AMDGPUMCInstLower.cpp
   AMDGPUMIRFormatter.cpp
   AMDGPUOpenCLEnqueuedBlockLowering.cpp
   AMDGPUPostLegalizerCombiner.cpp