1fe6060f1SDimitry Andric //===-- AMDGPULowerModuleLDSPass.cpp ------------------------------*- C++ -*-=//
2fe6060f1SDimitry Andric //
3fe6060f1SDimitry Andric // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4fe6060f1SDimitry Andric // See https://llvm.org/LICENSE.txt for license information.
5fe6060f1SDimitry Andric // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6fe6060f1SDimitry Andric //
7fe6060f1SDimitry Andric //===----------------------------------------------------------------------===//
8fe6060f1SDimitry Andric //
9bdd1243dSDimitry Andric // This pass eliminates local data store, LDS, uses from non-kernel functions.
10bdd1243dSDimitry Andric // LDS is contiguous memory allocated per kernel execution.
11fe6060f1SDimitry Andric //
12bdd1243dSDimitry Andric // Background.
13fe6060f1SDimitry Andric //
14bdd1243dSDimitry Andric // The programming model is global variables, or equivalently function local
15bdd1243dSDimitry Andric // static variables, accessible from kernels or other functions. For uses from
16bdd1243dSDimitry Andric // kernels this is straightforward - assign an integer to the kernel for the
17bdd1243dSDimitry Andric // memory required by all the variables combined, allocate them within that.
18bdd1243dSDimitry Andric // For uses from functions there are performance tradeoffs to choose between.
19bdd1243dSDimitry Andric //
20bdd1243dSDimitry Andric // This model means the GPU runtime can specify the amount of memory allocated.
21bdd1243dSDimitry Andric // If this is more than the kernel assumed, the excess can be made available
22bdd1243dSDimitry Andric // using a language specific feature, which IR represents as a variable with
2306c3fb27SDimitry Andric // no initializer. This feature is referred to here as "Dynamic LDS" and is
2406c3fb27SDimitry Andric // lowered slightly differently to the normal case.
25bdd1243dSDimitry Andric //
26bdd1243dSDimitry Andric // Consequences of this GPU feature:
27bdd1243dSDimitry Andric // - memory is limited and exceeding it halts compilation
28bdd1243dSDimitry Andric // - a global accessed by one kernel exists independent of other kernels
29bdd1243dSDimitry Andric // - a global exists independent of simultaneous execution of the same kernel
30bdd1243dSDimitry Andric // - the address of the global may be different from different kernels as they
31bdd1243dSDimitry Andric // do not alias, which permits only allocating variables they use
32bdd1243dSDimitry Andric // - if the address is allowed to differ, functions need help to find it
33bdd1243dSDimitry Andric //
34bdd1243dSDimitry Andric // Uses from kernels are implemented here by grouping them in a per-kernel
35bdd1243dSDimitry Andric // struct instance. This duplicates the variables, accurately modelling their
36bdd1243dSDimitry Andric // aliasing properties relative to a single global representation. It also
37bdd1243dSDimitry Andric // permits control over alignment via padding.
38bdd1243dSDimitry Andric //
39bdd1243dSDimitry Andric // Uses from functions are more complicated and the primary purpose of this
40bdd1243dSDimitry Andric // IR pass. Several different lowering are chosen between to meet requirements
41bdd1243dSDimitry Andric // to avoid allocating any LDS where it is not necessary, as that impacts
42bdd1243dSDimitry Andric // occupancy and may fail the compilation, while not imposing overhead on a
43bdd1243dSDimitry Andric // feature whose primary advantage over global memory is performance. The basic
44bdd1243dSDimitry Andric // design goal is to avoid one kernel imposing overhead on another.
45bdd1243dSDimitry Andric //
46bdd1243dSDimitry Andric // Implementation.
47bdd1243dSDimitry Andric //
48bdd1243dSDimitry Andric // LDS variables with constant annotation or non-undef initializer are passed
4981ad6265SDimitry Andric // through unchanged for simplification or error diagnostics in later passes.
50bdd1243dSDimitry Andric // Non-undef initializers are not yet implemented for LDS.
51fe6060f1SDimitry Andric //
52bdd1243dSDimitry Andric // LDS variables that are always allocated at the same address can be found
53bdd1243dSDimitry Andric // by lookup at that address. Otherwise runtime information/cost is required.
54fe6060f1SDimitry Andric //
55bdd1243dSDimitry Andric // The simplest strategy possible is to group all LDS variables in a single
56bdd1243dSDimitry Andric // struct and allocate that struct in every kernel such that the original
57bdd1243dSDimitry Andric // variables are always at the same address. LDS is however a limited resource
58bdd1243dSDimitry Andric // so this strategy is unusable in practice. It is not implemented here.
59bdd1243dSDimitry Andric //
60bdd1243dSDimitry Andric // Strategy | Precise allocation | Zero runtime cost | General purpose |
61bdd1243dSDimitry Andric // --------+--------------------+-------------------+-----------------+
62bdd1243dSDimitry Andric // Module | No | Yes | Yes |
63bdd1243dSDimitry Andric // Table | Yes | No | Yes |
64bdd1243dSDimitry Andric // Kernel | Yes | Yes | No |
65bdd1243dSDimitry Andric // Hybrid | Yes | Partial | Yes |
66bdd1243dSDimitry Andric //
6706c3fb27SDimitry Andric // "Module" spends LDS memory to save cycles. "Table" spends cycles and global
6806c3fb27SDimitry Andric // memory to save LDS. "Kernel" is as fast as kernel allocation but only works
6906c3fb27SDimitry Andric // for variables that are known reachable from a single kernel. "Hybrid" picks
7006c3fb27SDimitry Andric // between all three. When forced to choose between LDS and cycles we minimise
71bdd1243dSDimitry Andric // LDS use.
72bdd1243dSDimitry Andric
73bdd1243dSDimitry Andric // The "module" lowering implemented here finds LDS variables which are used by
74bdd1243dSDimitry Andric // non-kernel functions and creates a new struct with a field for each of those
75bdd1243dSDimitry Andric // LDS variables. Variables that are only used from kernels are excluded.
76bdd1243dSDimitry Andric //
77bdd1243dSDimitry Andric // The "table" lowering implemented here has three components.
78bdd1243dSDimitry Andric // First kernels are assigned a unique integer identifier which is available in
79bdd1243dSDimitry Andric // functions it calls through the intrinsic amdgcn_lds_kernel_id. The integer
80bdd1243dSDimitry Andric // is passed through a specific SGPR, thus works with indirect calls.
81bdd1243dSDimitry Andric // Second, each kernel allocates LDS variables independent of other kernels and
82bdd1243dSDimitry Andric // writes the addresses it chose for each variable into an array in consistent
83bdd1243dSDimitry Andric // order. If the kernel does not allocate a given variable, it writes undef to
84bdd1243dSDimitry Andric // the corresponding array location. These arrays are written to a constant
85bdd1243dSDimitry Andric // table in the order matching the kernel unique integer identifier.
86bdd1243dSDimitry Andric // Third, uses from non-kernel functions are replaced with a table lookup using
87bdd1243dSDimitry Andric // the intrinsic function to find the address of the variable.
88bdd1243dSDimitry Andric //
89bdd1243dSDimitry Andric // "Kernel" lowering is only applicable for variables that are unambiguously
90bdd1243dSDimitry Andric // reachable from exactly one kernel. For those cases, accesses to the variable
91bdd1243dSDimitry Andric // can be lowered to ConstantExpr address of a struct instance specific to that
92bdd1243dSDimitry Andric // one kernel. This is zero cost in space and in compute. It will raise a fatal
93bdd1243dSDimitry Andric // error on any variable that might be reachable from multiple kernels and is
94bdd1243dSDimitry Andric // thus most easily used as part of the hybrid lowering strategy.
95bdd1243dSDimitry Andric //
96bdd1243dSDimitry Andric // Hybrid lowering is a mixture of the above. It uses the zero cost kernel
97bdd1243dSDimitry Andric // lowering where it can. It lowers the variable accessed by the greatest
98bdd1243dSDimitry Andric // number of kernels using the module strategy as that is free for the first
99bdd1243dSDimitry Andric // variable. Any futher variables that can be lowered with the module strategy
100bdd1243dSDimitry Andric // without incurring LDS memory overhead are. The remaining ones are lowered
101bdd1243dSDimitry Andric // via table.
102bdd1243dSDimitry Andric //
103bdd1243dSDimitry Andric // Consequences
104bdd1243dSDimitry Andric // - No heuristics or user controlled magic numbers, hybrid is the right choice
105bdd1243dSDimitry Andric // - Kernels that don't use functions (or have had them all inlined) are not
106bdd1243dSDimitry Andric // affected by any lowering for kernels that do.
107bdd1243dSDimitry Andric // - Kernels that don't make indirect function calls are not affected by those
108bdd1243dSDimitry Andric // that do.
109bdd1243dSDimitry Andric // - Variables which are used by lots of kernels, e.g. those injected by a
110bdd1243dSDimitry Andric // language runtime in most kernels, are expected to have no overhead
111bdd1243dSDimitry Andric // - Implementations that instantiate templates per-kernel where those templates
112bdd1243dSDimitry Andric // use LDS are expected to hit the "Kernel" lowering strategy
113bdd1243dSDimitry Andric // - The runtime properties impose a cost in compiler implementation complexity
114fe6060f1SDimitry Andric //
11506c3fb27SDimitry Andric // Dynamic LDS implementation
11606c3fb27SDimitry Andric // Dynamic LDS is lowered similarly to the "table" strategy above and uses the
11706c3fb27SDimitry Andric // same intrinsic to identify which kernel is at the root of the dynamic call
11806c3fb27SDimitry Andric // graph. This relies on the specified behaviour that all dynamic LDS variables
11906c3fb27SDimitry Andric // alias one another, i.e. are at the same address, with respect to a given
12006c3fb27SDimitry Andric // kernel. Therefore this pass creates new dynamic LDS variables for each kernel
12106c3fb27SDimitry Andric // that allocates any dynamic LDS and builds a table of addresses out of those.
12206c3fb27SDimitry Andric // The AMDGPUPromoteAlloca pass skips kernels that use dynamic LDS.
12306c3fb27SDimitry Andric // The corresponding optimisation for "kernel" lowering where the table lookup
12406c3fb27SDimitry Andric // is elided is not implemented.
12506c3fb27SDimitry Andric //
12606c3fb27SDimitry Andric //
12706c3fb27SDimitry Andric // Implementation notes / limitations
12806c3fb27SDimitry Andric // A single LDS global variable represents an instance per kernel that can reach
12906c3fb27SDimitry Andric // said variables. This pass essentially specialises said variables per kernel.
13006c3fb27SDimitry Andric // Handling ConstantExpr during the pass complicated this significantly so now
13106c3fb27SDimitry Andric // all ConstantExpr uses of LDS variables are expanded to instructions. This
13206c3fb27SDimitry Andric // may need amending when implementing non-undef initialisers.
13306c3fb27SDimitry Andric //
13406c3fb27SDimitry Andric // Lowering is split between this IR pass and the back end. This pass chooses
13506c3fb27SDimitry Andric // where given variables should be allocated and marks them with metadata,
13606c3fb27SDimitry Andric // MD_absolute_symbol. The backend places the variables in coincidentally the
13706c3fb27SDimitry Andric // same location and raises a fatal error if something has gone awry. This works
13806c3fb27SDimitry Andric // in practice because the only pass between this one and the backend that
13906c3fb27SDimitry Andric // changes LDS is PromoteAlloca and the changes it makes do not conflict.
14006c3fb27SDimitry Andric //
14106c3fb27SDimitry Andric // Addresses are written to constant global arrays based on the same metadata.
14206c3fb27SDimitry Andric //
14306c3fb27SDimitry Andric // The backend lowers LDS variables in the order of traversal of the function.
14406c3fb27SDimitry Andric // This is at odds with the deterministic layout required. The workaround is to
14506c3fb27SDimitry Andric // allocate the fixed-address variables immediately upon starting the function
14606c3fb27SDimitry Andric // where they can be placed as intended. This requires a means of mapping from
14706c3fb27SDimitry Andric // the function to the variables that it allocates. For the module scope lds,
14806c3fb27SDimitry Andric // this is via metadata indicating whether the variable is not required. If a
14906c3fb27SDimitry Andric // pass deletes that metadata, a fatal error on disagreement with the absolute
15006c3fb27SDimitry Andric // symbol metadata will occur. For kernel scope and dynamic, this is by _name_
15106c3fb27SDimitry Andric // correspondence between the function and the variable. It requires the
15206c3fb27SDimitry Andric // kernel to have a name (which is only a limitation for tests in practice) and
15306c3fb27SDimitry Andric // for nothing to rename the corresponding symbols. This is a hazard if the pass
15406c3fb27SDimitry Andric // is run multiple times during debugging. Alternative schemes considered all
15506c3fb27SDimitry Andric // involve bespoke metadata.
15606c3fb27SDimitry Andric //
15706c3fb27SDimitry Andric // If the name correspondence can be replaced, multiple distinct kernels that
15806c3fb27SDimitry Andric // have the same memory layout can map to the same kernel id (as the address
15906c3fb27SDimitry Andric // itself is handled by the absolute symbol metadata) and that will allow more
16006c3fb27SDimitry Andric // uses of the "kernel" style faster lowering and reduce the size of the lookup
16106c3fb27SDimitry Andric // tables.
16206c3fb27SDimitry Andric //
16306c3fb27SDimitry Andric // There is a test that checks this does not fire for a graphics shader. This
16406c3fb27SDimitry Andric // lowering is expected to work for graphics if the isKernel test is changed.
16506c3fb27SDimitry Andric //
16606c3fb27SDimitry Andric // The current markUsedByKernel is sufficient for PromoteAlloca but is elided
16706c3fb27SDimitry Andric // before codegen. Replacing this with an equivalent intrinsic which lasts until
16806c3fb27SDimitry Andric // shortly after the machine function lowering of LDS would help break the name
16906c3fb27SDimitry Andric // mapping. The other part needed is probably to amend PromoteAlloca to embed
17006c3fb27SDimitry Andric // the LDS variables it creates in the same struct created here. That avoids the
17106c3fb27SDimitry Andric // current hazard where a PromoteAlloca LDS variable might be allocated before
17206c3fb27SDimitry Andric // the kernel scope (and thus error on the address check). Given a new invariant
17306c3fb27SDimitry Andric // that no LDS variables exist outside of the structs managed here, and an
17406c3fb27SDimitry Andric // intrinsic that lasts until after the LDS frame lowering, it should be
17506c3fb27SDimitry Andric // possible to drop the name mapping and fold equivalent memory layouts.
17606c3fb27SDimitry Andric //
177fe6060f1SDimitry Andric //===----------------------------------------------------------------------===//
178fe6060f1SDimitry Andric
179fe6060f1SDimitry Andric #include "AMDGPU.h"
1805f757f3fSDimitry Andric #include "AMDGPUTargetMachine.h"
181fe6060f1SDimitry Andric #include "Utils/AMDGPUBaseInfo.h"
18281ad6265SDimitry Andric #include "Utils/AMDGPUMemoryUtils.h"
183972a253aSDimitry Andric #include "llvm/ADT/BitVector.h"
184972a253aSDimitry Andric #include "llvm/ADT/DenseMap.h"
185bdd1243dSDimitry Andric #include "llvm/ADT/DenseSet.h"
186fe6060f1SDimitry Andric #include "llvm/ADT/STLExtras.h"
187bdd1243dSDimitry Andric #include "llvm/ADT/SetOperations.h"
18881ad6265SDimitry Andric #include "llvm/Analysis/CallGraph.h"
1895f757f3fSDimitry Andric #include "llvm/CodeGen/TargetPassConfig.h"
190fe6060f1SDimitry Andric #include "llvm/IR/Constants.h"
191fe6060f1SDimitry Andric #include "llvm/IR/DerivedTypes.h"
192fe6060f1SDimitry Andric #include "llvm/IR/IRBuilder.h"
193fe6060f1SDimitry Andric #include "llvm/IR/InlineAsm.h"
194fe6060f1SDimitry Andric #include "llvm/IR/Instructions.h"
195bdd1243dSDimitry Andric #include "llvm/IR/IntrinsicsAMDGPU.h"
196349cc55cSDimitry Andric #include "llvm/IR/MDBuilder.h"
19706c3fb27SDimitry Andric #include "llvm/IR/ReplaceConstant.h"
198fe6060f1SDimitry Andric #include "llvm/InitializePasses.h"
199fe6060f1SDimitry Andric #include "llvm/Pass.h"
200fe6060f1SDimitry Andric #include "llvm/Support/CommandLine.h"
201fe6060f1SDimitry Andric #include "llvm/Support/Debug.h"
20206c3fb27SDimitry Andric #include "llvm/Support/Format.h"
203fe6060f1SDimitry Andric #include "llvm/Support/OptimizedStructLayout.h"
20406c3fb27SDimitry Andric #include "llvm/Support/raw_ostream.h"
205bdd1243dSDimitry Andric #include "llvm/Transforms/Utils/BasicBlockUtils.h"
206fe6060f1SDimitry Andric #include "llvm/Transforms/Utils/ModuleUtils.h"
207bdd1243dSDimitry Andric
208fe6060f1SDimitry Andric #include <vector>
209fe6060f1SDimitry Andric
210bdd1243dSDimitry Andric #include <cstdio>
211bdd1243dSDimitry Andric
212fe6060f1SDimitry Andric #define DEBUG_TYPE "amdgpu-lower-module-lds"
213fe6060f1SDimitry Andric
214fe6060f1SDimitry Andric using namespace llvm;
215fe6060f1SDimitry Andric
216bdd1243dSDimitry Andric namespace {
217bdd1243dSDimitry Andric
218bdd1243dSDimitry Andric cl::opt<bool> SuperAlignLDSGlobals(
219fe6060f1SDimitry Andric "amdgpu-super-align-lds-globals",
220fe6060f1SDimitry Andric cl::desc("Increase alignment of LDS if it is not on align boundary"),
221fe6060f1SDimitry Andric cl::init(true), cl::Hidden);
222fe6060f1SDimitry Andric
223bdd1243dSDimitry Andric enum class LoweringKind { module, table, kernel, hybrid };
224bdd1243dSDimitry Andric cl::opt<LoweringKind> LoweringKindLoc(
225bdd1243dSDimitry Andric "amdgpu-lower-module-lds-strategy",
226bdd1243dSDimitry Andric cl::desc("Specify lowering strategy for function LDS access:"), cl::Hidden,
22706c3fb27SDimitry Andric cl::init(LoweringKind::hybrid),
228bdd1243dSDimitry Andric cl::values(
229bdd1243dSDimitry Andric clEnumValN(LoweringKind::table, "table", "Lower via table lookup"),
230bdd1243dSDimitry Andric clEnumValN(LoweringKind::module, "module", "Lower via module struct"),
231bdd1243dSDimitry Andric clEnumValN(
232bdd1243dSDimitry Andric LoweringKind::kernel, "kernel",
233bdd1243dSDimitry Andric "Lower variables reachable from one kernel, otherwise abort"),
234bdd1243dSDimitry Andric clEnumValN(LoweringKind::hybrid, "hybrid",
235bdd1243dSDimitry Andric "Lower via mixture of above strategies")));
236bdd1243dSDimitry Andric
isKernelLDS(const Function * F)237bdd1243dSDimitry Andric bool isKernelLDS(const Function *F) {
238bdd1243dSDimitry Andric // Some weirdness here. AMDGPU::isKernelCC does not call into
239bdd1243dSDimitry Andric // AMDGPU::isKernel with the calling conv, it instead calls into
240bdd1243dSDimitry Andric // isModuleEntryFunction which returns true for more calling conventions
241bdd1243dSDimitry Andric // than AMDGPU::isKernel does. There's a FIXME on AMDGPU::isKernel.
242bdd1243dSDimitry Andric // There's also a test that checks that the LDS lowering does not hit on
243bdd1243dSDimitry Andric // a graphics shader, denoted amdgpu_ps, so stay with the limited case.
244bdd1243dSDimitry Andric // Putting LDS in the name of the function to draw attention to this.
245bdd1243dSDimitry Andric return AMDGPU::isKernel(F->getCallingConv());
246bdd1243dSDimitry Andric }
247bdd1243dSDimitry Andric
sortByName(std::vector<T> && V)24806c3fb27SDimitry Andric template <typename T> std::vector<T> sortByName(std::vector<T> &&V) {
24906c3fb27SDimitry Andric llvm::sort(V.begin(), V.end(), [](const auto *L, const auto *R) {
25006c3fb27SDimitry Andric return L->getName() < R->getName();
25106c3fb27SDimitry Andric });
25206c3fb27SDimitry Andric return {std::move(V)};
25306c3fb27SDimitry Andric }
25406c3fb27SDimitry Andric
2555f757f3fSDimitry Andric class AMDGPULowerModuleLDS {
2565f757f3fSDimitry Andric const AMDGPUTargetMachine &TM;
257fe6060f1SDimitry Andric
258fe6060f1SDimitry Andric static void
removeLocalVarsFromUsedLists(Module & M,const DenseSet<GlobalVariable * > & LocalVars)259bdd1243dSDimitry Andric removeLocalVarsFromUsedLists(Module &M,
260bdd1243dSDimitry Andric const DenseSet<GlobalVariable *> &LocalVars) {
261972a253aSDimitry Andric // The verifier rejects used lists containing an inttoptr of a constant
262972a253aSDimitry Andric // so remove the variables from these lists before replaceAllUsesWith
263bdd1243dSDimitry Andric SmallPtrSet<Constant *, 8> LocalVarsSet;
2640eae32dcSDimitry Andric for (GlobalVariable *LocalVar : LocalVars)
265bdd1243dSDimitry Andric LocalVarsSet.insert(cast<Constant>(LocalVar->stripPointerCasts()));
266bdd1243dSDimitry Andric
267bdd1243dSDimitry Andric removeFromUsedLists(
268bdd1243dSDimitry Andric M, [&LocalVarsSet](Constant *C) { return LocalVarsSet.count(C); });
269bdd1243dSDimitry Andric
270bdd1243dSDimitry Andric for (GlobalVariable *LocalVar : LocalVars)
271bdd1243dSDimitry Andric LocalVar->removeDeadConstantUsers();
272fe6060f1SDimitry Andric }
273fe6060f1SDimitry Andric
markUsedByKernel(Function * Func,GlobalVariable * SGV)27406c3fb27SDimitry Andric static void markUsedByKernel(Function *Func, GlobalVariable *SGV) {
275fe6060f1SDimitry Andric // The llvm.amdgcn.module.lds instance is implicitly used by all kernels
276fe6060f1SDimitry Andric // that might call a function which accesses a field within it. This is
277fe6060f1SDimitry Andric // presently approximated to 'all kernels' if there are any such functions
278349cc55cSDimitry Andric // in the module. This implicit use is redefined as an explicit use here so
279fe6060f1SDimitry Andric // that later passes, specifically PromoteAlloca, account for the required
280fe6060f1SDimitry Andric // memory without any knowledge of this transform.
281fe6060f1SDimitry Andric
282fe6060f1SDimitry Andric // An operand bundle on llvm.donothing works because the call instruction
283fe6060f1SDimitry Andric // survives until after the last pass that needs to account for LDS. It is
284fe6060f1SDimitry Andric // better than inline asm as the latter survives until the end of codegen. A
285fe6060f1SDimitry Andric // totally robust solution would be a function with the same semantics as
286fe6060f1SDimitry Andric // llvm.donothing that takes a pointer to the instance and is lowered to a
287fe6060f1SDimitry Andric // no-op after LDS is allocated, but that is not presently necessary.
288fe6060f1SDimitry Andric
28906c3fb27SDimitry Andric // This intrinsic is eliminated shortly before instruction selection. It
29006c3fb27SDimitry Andric // does not suffice to indicate to ISel that a given global which is not
29106c3fb27SDimitry Andric // immediately used by the kernel must still be allocated by it. An
29206c3fb27SDimitry Andric // equivalent target specific intrinsic which lasts until immediately after
29306c3fb27SDimitry Andric // codegen would suffice for that, but one would still need to ensure that
29406c3fb27SDimitry Andric // the variables are allocated in the anticpated order.
2955f757f3fSDimitry Andric BasicBlock *Entry = &Func->getEntryBlock();
2965f757f3fSDimitry Andric IRBuilder<> Builder(Entry, Entry->getFirstNonPHIIt());
297fe6060f1SDimitry Andric
298fe6060f1SDimitry Andric Function *Decl =
299fe6060f1SDimitry Andric Intrinsic::getDeclaration(Func->getParent(), Intrinsic::donothing, {});
300fe6060f1SDimitry Andric
30106c3fb27SDimitry Andric Value *UseInstance[1] = {
30206c3fb27SDimitry Andric Builder.CreateConstInBoundsGEP1_32(SGV->getValueType(), SGV, 0)};
303fe6060f1SDimitry Andric
30406c3fb27SDimitry Andric Builder.CreateCall(
30506c3fb27SDimitry Andric Decl, {}, {OperandBundleDefT<Value *>("ExplicitUse", UseInstance)});
306fe6060f1SDimitry Andric }
307fe6060f1SDimitry Andric
eliminateConstantExprUsesOfLDSFromAllInstructions(Module & M)308bdd1243dSDimitry Andric static bool eliminateConstantExprUsesOfLDSFromAllInstructions(Module &M) {
309bdd1243dSDimitry Andric // Constants are uniqued within LLVM. A ConstantExpr referring to a LDS
310bdd1243dSDimitry Andric // global may have uses from multiple different functions as a result.
311bdd1243dSDimitry Andric // This pass specialises LDS variables with respect to the kernel that
312bdd1243dSDimitry Andric // allocates them.
313bdd1243dSDimitry Andric
31406c3fb27SDimitry Andric // This is semantically equivalent to (the unimplemented as slow):
315bdd1243dSDimitry Andric // for (auto &F : M.functions())
316bdd1243dSDimitry Andric // for (auto &BB : F)
317bdd1243dSDimitry Andric // for (auto &I : BB)
318bdd1243dSDimitry Andric // for (Use &Op : I.operands())
319bdd1243dSDimitry Andric // if (constantExprUsesLDS(Op))
320bdd1243dSDimitry Andric // replaceConstantExprInFunction(I, Op);
321bdd1243dSDimitry Andric
32206c3fb27SDimitry Andric SmallVector<Constant *> LDSGlobals;
323bdd1243dSDimitry Andric for (auto &GV : M.globals())
324bdd1243dSDimitry Andric if (AMDGPU::isLDSVariableToLower(GV))
32506c3fb27SDimitry Andric LDSGlobals.push_back(&GV);
326bdd1243dSDimitry Andric
32706c3fb27SDimitry Andric return convertUsersOfConstantsToInstructions(LDSGlobals);
328bdd1243dSDimitry Andric }
329bdd1243dSDimitry Andric
330fe6060f1SDimitry Andric public:
AMDGPULowerModuleLDS(const AMDGPUTargetMachine & TM_)3315f757f3fSDimitry Andric AMDGPULowerModuleLDS(const AMDGPUTargetMachine &TM_) : TM(TM_) {}
332fe6060f1SDimitry Andric
333bdd1243dSDimitry Andric using FunctionVariableMap = DenseMap<Function *, DenseSet<GlobalVariable *>>;
334bdd1243dSDimitry Andric
335bdd1243dSDimitry Andric using VariableFunctionMap = DenseMap<GlobalVariable *, DenseSet<Function *>>;
336bdd1243dSDimitry Andric
getUsesOfLDSByFunction(CallGraph const & CG,Module & M,FunctionVariableMap & kernels,FunctionVariableMap & functions)337bdd1243dSDimitry Andric static void getUsesOfLDSByFunction(CallGraph const &CG, Module &M,
338bdd1243dSDimitry Andric FunctionVariableMap &kernels,
339bdd1243dSDimitry Andric FunctionVariableMap &functions) {
340bdd1243dSDimitry Andric
341bdd1243dSDimitry Andric // Get uses from the current function, excluding uses by called functions
342bdd1243dSDimitry Andric // Two output variables to avoid walking the globals list twice
343bdd1243dSDimitry Andric for (auto &GV : M.globals()) {
344bdd1243dSDimitry Andric if (!AMDGPU::isLDSVariableToLower(GV)) {
345bdd1243dSDimitry Andric continue;
346bdd1243dSDimitry Andric }
347bdd1243dSDimitry Andric
34806c3fb27SDimitry Andric if (GV.isAbsoluteSymbolRef()) {
34906c3fb27SDimitry Andric report_fatal_error(
35006c3fb27SDimitry Andric "LDS variables with absolute addresses are unimplemented.");
35106c3fb27SDimitry Andric }
35206c3fb27SDimitry Andric
353bdd1243dSDimitry Andric for (User *V : GV.users()) {
354bdd1243dSDimitry Andric if (auto *I = dyn_cast<Instruction>(V)) {
355bdd1243dSDimitry Andric Function *F = I->getFunction();
356bdd1243dSDimitry Andric if (isKernelLDS(F)) {
357bdd1243dSDimitry Andric kernels[F].insert(&GV);
358bdd1243dSDimitry Andric } else {
359bdd1243dSDimitry Andric functions[F].insert(&GV);
360bdd1243dSDimitry Andric }
361bdd1243dSDimitry Andric }
362bdd1243dSDimitry Andric }
363bdd1243dSDimitry Andric }
364bdd1243dSDimitry Andric }
365bdd1243dSDimitry Andric
366bdd1243dSDimitry Andric struct LDSUsesInfoTy {
367bdd1243dSDimitry Andric FunctionVariableMap direct_access;
368bdd1243dSDimitry Andric FunctionVariableMap indirect_access;
369bdd1243dSDimitry Andric };
370bdd1243dSDimitry Andric
getTransitiveUsesOfLDS(CallGraph const & CG,Module & M)371bdd1243dSDimitry Andric static LDSUsesInfoTy getTransitiveUsesOfLDS(CallGraph const &CG, Module &M) {
372bdd1243dSDimitry Andric
373bdd1243dSDimitry Andric FunctionVariableMap direct_map_kernel;
374bdd1243dSDimitry Andric FunctionVariableMap direct_map_function;
375bdd1243dSDimitry Andric getUsesOfLDSByFunction(CG, M, direct_map_kernel, direct_map_function);
376bdd1243dSDimitry Andric
377bdd1243dSDimitry Andric // Collect variables that are used by functions whose address has escaped
378bdd1243dSDimitry Andric DenseSet<GlobalVariable *> VariablesReachableThroughFunctionPointer;
379bdd1243dSDimitry Andric for (Function &F : M.functions()) {
380bdd1243dSDimitry Andric if (!isKernelLDS(&F))
381bdd1243dSDimitry Andric if (F.hasAddressTaken(nullptr,
382bdd1243dSDimitry Andric /* IgnoreCallbackUses */ false,
383bdd1243dSDimitry Andric /* IgnoreAssumeLikeCalls */ false,
384bdd1243dSDimitry Andric /* IgnoreLLVMUsed */ true,
385bdd1243dSDimitry Andric /* IgnoreArcAttachedCall */ false)) {
386bdd1243dSDimitry Andric set_union(VariablesReachableThroughFunctionPointer,
387bdd1243dSDimitry Andric direct_map_function[&F]);
388bdd1243dSDimitry Andric }
389bdd1243dSDimitry Andric }
390bdd1243dSDimitry Andric
391bdd1243dSDimitry Andric auto functionMakesUnknownCall = [&](const Function *F) -> bool {
392bdd1243dSDimitry Andric assert(!F->isDeclaration());
39306c3fb27SDimitry Andric for (const CallGraphNode::CallRecord &R : *CG[F]) {
394bdd1243dSDimitry Andric if (!R.second->getFunction()) {
395bdd1243dSDimitry Andric return true;
396bdd1243dSDimitry Andric }
397bdd1243dSDimitry Andric }
398bdd1243dSDimitry Andric return false;
399bdd1243dSDimitry Andric };
400bdd1243dSDimitry Andric
401bdd1243dSDimitry Andric // Work out which variables are reachable through function calls
402bdd1243dSDimitry Andric FunctionVariableMap transitive_map_function = direct_map_function;
403bdd1243dSDimitry Andric
404bdd1243dSDimitry Andric // If the function makes any unknown call, assume the worst case that it can
405bdd1243dSDimitry Andric // access all variables accessed by functions whose address escaped
406bdd1243dSDimitry Andric for (Function &F : M.functions()) {
407bdd1243dSDimitry Andric if (!F.isDeclaration() && functionMakesUnknownCall(&F)) {
408bdd1243dSDimitry Andric if (!isKernelLDS(&F)) {
409bdd1243dSDimitry Andric set_union(transitive_map_function[&F],
410bdd1243dSDimitry Andric VariablesReachableThroughFunctionPointer);
411bdd1243dSDimitry Andric }
412bdd1243dSDimitry Andric }
413bdd1243dSDimitry Andric }
414bdd1243dSDimitry Andric
415bdd1243dSDimitry Andric // Direct implementation of collecting all variables reachable from each
416bdd1243dSDimitry Andric // function
417bdd1243dSDimitry Andric for (Function &Func : M.functions()) {
418bdd1243dSDimitry Andric if (Func.isDeclaration() || isKernelLDS(&Func))
419bdd1243dSDimitry Andric continue;
420bdd1243dSDimitry Andric
421bdd1243dSDimitry Andric DenseSet<Function *> seen; // catches cycles
422bdd1243dSDimitry Andric SmallVector<Function *, 4> wip{&Func};
423bdd1243dSDimitry Andric
424bdd1243dSDimitry Andric while (!wip.empty()) {
425bdd1243dSDimitry Andric Function *F = wip.pop_back_val();
426bdd1243dSDimitry Andric
427bdd1243dSDimitry Andric // Can accelerate this by referring to transitive map for functions that
428bdd1243dSDimitry Andric // have already been computed, with more care than this
429bdd1243dSDimitry Andric set_union(transitive_map_function[&Func], direct_map_function[F]);
430bdd1243dSDimitry Andric
43106c3fb27SDimitry Andric for (const CallGraphNode::CallRecord &R : *CG[F]) {
432bdd1243dSDimitry Andric Function *ith = R.second->getFunction();
433bdd1243dSDimitry Andric if (ith) {
434bdd1243dSDimitry Andric if (!seen.contains(ith)) {
435bdd1243dSDimitry Andric seen.insert(ith);
436bdd1243dSDimitry Andric wip.push_back(ith);
437bdd1243dSDimitry Andric }
438bdd1243dSDimitry Andric }
439bdd1243dSDimitry Andric }
440bdd1243dSDimitry Andric }
441bdd1243dSDimitry Andric }
442bdd1243dSDimitry Andric
443bdd1243dSDimitry Andric // direct_map_kernel lists which variables are used by the kernel
444bdd1243dSDimitry Andric // find the variables which are used through a function call
445bdd1243dSDimitry Andric FunctionVariableMap indirect_map_kernel;
446bdd1243dSDimitry Andric
447bdd1243dSDimitry Andric for (Function &Func : M.functions()) {
448bdd1243dSDimitry Andric if (Func.isDeclaration() || !isKernelLDS(&Func))
449bdd1243dSDimitry Andric continue;
450bdd1243dSDimitry Andric
45106c3fb27SDimitry Andric for (const CallGraphNode::CallRecord &R : *CG[&Func]) {
452bdd1243dSDimitry Andric Function *ith = R.second->getFunction();
453bdd1243dSDimitry Andric if (ith) {
454bdd1243dSDimitry Andric set_union(indirect_map_kernel[&Func], transitive_map_function[ith]);
455bdd1243dSDimitry Andric } else {
456bdd1243dSDimitry Andric set_union(indirect_map_kernel[&Func],
457bdd1243dSDimitry Andric VariablesReachableThroughFunctionPointer);
458bdd1243dSDimitry Andric }
459bdd1243dSDimitry Andric }
460bdd1243dSDimitry Andric }
461bdd1243dSDimitry Andric
462bdd1243dSDimitry Andric return {std::move(direct_map_kernel), std::move(indirect_map_kernel)};
463bdd1243dSDimitry Andric }
464bdd1243dSDimitry Andric
465bdd1243dSDimitry Andric struct LDSVariableReplacement {
466bdd1243dSDimitry Andric GlobalVariable *SGV = nullptr;
467bdd1243dSDimitry Andric DenseMap<GlobalVariable *, Constant *> LDSVarsToConstantGEP;
468bdd1243dSDimitry Andric };
469bdd1243dSDimitry Andric
470bdd1243dSDimitry Andric // remap from lds global to a constantexpr gep to where it has been moved to
471bdd1243dSDimitry Andric // for each kernel
472bdd1243dSDimitry Andric // an array with an element for each kernel containing where the corresponding
473bdd1243dSDimitry Andric // variable was remapped to
474bdd1243dSDimitry Andric
getAddressesOfVariablesInKernel(LLVMContext & Ctx,ArrayRef<GlobalVariable * > Variables,const DenseMap<GlobalVariable *,Constant * > & LDSVarsToConstantGEP)475bdd1243dSDimitry Andric static Constant *getAddressesOfVariablesInKernel(
476bdd1243dSDimitry Andric LLVMContext &Ctx, ArrayRef<GlobalVariable *> Variables,
47706c3fb27SDimitry Andric const DenseMap<GlobalVariable *, Constant *> &LDSVarsToConstantGEP) {
478bdd1243dSDimitry Andric // Create a ConstantArray containing the address of each Variable within the
479bdd1243dSDimitry Andric // kernel corresponding to LDSVarsToConstantGEP, or poison if that kernel
480bdd1243dSDimitry Andric // does not allocate it
481bdd1243dSDimitry Andric // TODO: Drop the ptrtoint conversion
482bdd1243dSDimitry Andric
483bdd1243dSDimitry Andric Type *I32 = Type::getInt32Ty(Ctx);
484bdd1243dSDimitry Andric
485bdd1243dSDimitry Andric ArrayType *KernelOffsetsType = ArrayType::get(I32, Variables.size());
486bdd1243dSDimitry Andric
487bdd1243dSDimitry Andric SmallVector<Constant *> Elements;
488bdd1243dSDimitry Andric for (size_t i = 0; i < Variables.size(); i++) {
489bdd1243dSDimitry Andric GlobalVariable *GV = Variables[i];
49006c3fb27SDimitry Andric auto ConstantGepIt = LDSVarsToConstantGEP.find(GV);
49106c3fb27SDimitry Andric if (ConstantGepIt != LDSVarsToConstantGEP.end()) {
49206c3fb27SDimitry Andric auto elt = ConstantExpr::getPtrToInt(ConstantGepIt->second, I32);
493bdd1243dSDimitry Andric Elements.push_back(elt);
494bdd1243dSDimitry Andric } else {
495bdd1243dSDimitry Andric Elements.push_back(PoisonValue::get(I32));
496bdd1243dSDimitry Andric }
497bdd1243dSDimitry Andric }
498bdd1243dSDimitry Andric return ConstantArray::get(KernelOffsetsType, Elements);
499bdd1243dSDimitry Andric }
500bdd1243dSDimitry Andric
buildLookupTable(Module & M,ArrayRef<GlobalVariable * > Variables,ArrayRef<Function * > kernels,DenseMap<Function *,LDSVariableReplacement> & KernelToReplacement)501bdd1243dSDimitry Andric static GlobalVariable *buildLookupTable(
502bdd1243dSDimitry Andric Module &M, ArrayRef<GlobalVariable *> Variables,
503bdd1243dSDimitry Andric ArrayRef<Function *> kernels,
504bdd1243dSDimitry Andric DenseMap<Function *, LDSVariableReplacement> &KernelToReplacement) {
505bdd1243dSDimitry Andric if (Variables.empty()) {
506bdd1243dSDimitry Andric return nullptr;
507bdd1243dSDimitry Andric }
508bdd1243dSDimitry Andric LLVMContext &Ctx = M.getContext();
509bdd1243dSDimitry Andric
510bdd1243dSDimitry Andric const size_t NumberVariables = Variables.size();
511bdd1243dSDimitry Andric const size_t NumberKernels = kernels.size();
512bdd1243dSDimitry Andric
513bdd1243dSDimitry Andric ArrayType *KernelOffsetsType =
514bdd1243dSDimitry Andric ArrayType::get(Type::getInt32Ty(Ctx), NumberVariables);
515bdd1243dSDimitry Andric
516bdd1243dSDimitry Andric ArrayType *AllKernelsOffsetsType =
517bdd1243dSDimitry Andric ArrayType::get(KernelOffsetsType, NumberKernels);
518bdd1243dSDimitry Andric
51906c3fb27SDimitry Andric Constant *Missing = PoisonValue::get(KernelOffsetsType);
520bdd1243dSDimitry Andric std::vector<Constant *> overallConstantExprElts(NumberKernels);
521bdd1243dSDimitry Andric for (size_t i = 0; i < NumberKernels; i++) {
52206c3fb27SDimitry Andric auto Replacement = KernelToReplacement.find(kernels[i]);
52306c3fb27SDimitry Andric overallConstantExprElts[i] =
52406c3fb27SDimitry Andric (Replacement == KernelToReplacement.end())
52506c3fb27SDimitry Andric ? Missing
52606c3fb27SDimitry Andric : getAddressesOfVariablesInKernel(
52706c3fb27SDimitry Andric Ctx, Variables, Replacement->second.LDSVarsToConstantGEP);
528bdd1243dSDimitry Andric }
529bdd1243dSDimitry Andric
530bdd1243dSDimitry Andric Constant *init =
531bdd1243dSDimitry Andric ConstantArray::get(AllKernelsOffsetsType, overallConstantExprElts);
532bdd1243dSDimitry Andric
533bdd1243dSDimitry Andric return new GlobalVariable(
534bdd1243dSDimitry Andric M, AllKernelsOffsetsType, true, GlobalValue::InternalLinkage, init,
535bdd1243dSDimitry Andric "llvm.amdgcn.lds.offset.table", nullptr, GlobalValue::NotThreadLocal,
536bdd1243dSDimitry Andric AMDGPUAS::CONSTANT_ADDRESS);
537bdd1243dSDimitry Andric }
538bdd1243dSDimitry Andric
replaceUseWithTableLookup(Module & M,IRBuilder<> & Builder,GlobalVariable * LookupTable,GlobalVariable * GV,Use & U,Value * OptionalIndex)53906c3fb27SDimitry Andric void replaceUseWithTableLookup(Module &M, IRBuilder<> &Builder,
54006c3fb27SDimitry Andric GlobalVariable *LookupTable,
54106c3fb27SDimitry Andric GlobalVariable *GV, Use &U,
54206c3fb27SDimitry Andric Value *OptionalIndex) {
54306c3fb27SDimitry Andric // Table is a constant array of the same length as OrderedKernels
544bdd1243dSDimitry Andric LLVMContext &Ctx = M.getContext();
545bdd1243dSDimitry Andric Type *I32 = Type::getInt32Ty(Ctx);
54606c3fb27SDimitry Andric auto *I = cast<Instruction>(U.getUser());
547bdd1243dSDimitry Andric
54806c3fb27SDimitry Andric Value *tableKernelIndex = getTableLookupKernelIndex(M, I->getFunction());
549bdd1243dSDimitry Andric
550bdd1243dSDimitry Andric if (auto *Phi = dyn_cast<PHINode>(I)) {
551bdd1243dSDimitry Andric BasicBlock *BB = Phi->getIncomingBlock(U);
552bdd1243dSDimitry Andric Builder.SetInsertPoint(&(*(BB->getFirstInsertionPt())));
553bdd1243dSDimitry Andric } else {
554bdd1243dSDimitry Andric Builder.SetInsertPoint(I);
555bdd1243dSDimitry Andric }
556bdd1243dSDimitry Andric
55706c3fb27SDimitry Andric SmallVector<Value *, 3> GEPIdx = {
558bdd1243dSDimitry Andric ConstantInt::get(I32, 0),
559bdd1243dSDimitry Andric tableKernelIndex,
560bdd1243dSDimitry Andric };
56106c3fb27SDimitry Andric if (OptionalIndex)
56206c3fb27SDimitry Andric GEPIdx.push_back(OptionalIndex);
563bdd1243dSDimitry Andric
564bdd1243dSDimitry Andric Value *Address = Builder.CreateInBoundsGEP(
565bdd1243dSDimitry Andric LookupTable->getValueType(), LookupTable, GEPIdx, GV->getName());
566bdd1243dSDimitry Andric
567bdd1243dSDimitry Andric Value *loaded = Builder.CreateLoad(I32, Address);
568bdd1243dSDimitry Andric
569bdd1243dSDimitry Andric Value *replacement =
570bdd1243dSDimitry Andric Builder.CreateIntToPtr(loaded, GV->getType(), GV->getName());
571bdd1243dSDimitry Andric
572bdd1243dSDimitry Andric U.set(replacement);
573bdd1243dSDimitry Andric }
57406c3fb27SDimitry Andric
replaceUsesInInstructionsWithTableLookup(Module & M,ArrayRef<GlobalVariable * > ModuleScopeVariables,GlobalVariable * LookupTable)57506c3fb27SDimitry Andric void replaceUsesInInstructionsWithTableLookup(
57606c3fb27SDimitry Andric Module &M, ArrayRef<GlobalVariable *> ModuleScopeVariables,
57706c3fb27SDimitry Andric GlobalVariable *LookupTable) {
57806c3fb27SDimitry Andric
57906c3fb27SDimitry Andric LLVMContext &Ctx = M.getContext();
58006c3fb27SDimitry Andric IRBuilder<> Builder(Ctx);
58106c3fb27SDimitry Andric Type *I32 = Type::getInt32Ty(Ctx);
58206c3fb27SDimitry Andric
58306c3fb27SDimitry Andric for (size_t Index = 0; Index < ModuleScopeVariables.size(); Index++) {
58406c3fb27SDimitry Andric auto *GV = ModuleScopeVariables[Index];
58506c3fb27SDimitry Andric
58606c3fb27SDimitry Andric for (Use &U : make_early_inc_range(GV->uses())) {
58706c3fb27SDimitry Andric auto *I = dyn_cast<Instruction>(U.getUser());
58806c3fb27SDimitry Andric if (!I)
58906c3fb27SDimitry Andric continue;
59006c3fb27SDimitry Andric
59106c3fb27SDimitry Andric replaceUseWithTableLookup(M, Builder, LookupTable, GV, U,
59206c3fb27SDimitry Andric ConstantInt::get(I32, Index));
59306c3fb27SDimitry Andric }
594bdd1243dSDimitry Andric }
595bdd1243dSDimitry Andric }
596bdd1243dSDimitry Andric
kernelsThatIndirectlyAccessAnyOfPassedVariables(Module & M,LDSUsesInfoTy & LDSUsesInfo,DenseSet<GlobalVariable * > const & VariableSet)597bdd1243dSDimitry Andric static DenseSet<Function *> kernelsThatIndirectlyAccessAnyOfPassedVariables(
598bdd1243dSDimitry Andric Module &M, LDSUsesInfoTy &LDSUsesInfo,
599bdd1243dSDimitry Andric DenseSet<GlobalVariable *> const &VariableSet) {
600bdd1243dSDimitry Andric
601bdd1243dSDimitry Andric DenseSet<Function *> KernelSet;
602bdd1243dSDimitry Andric
60306c3fb27SDimitry Andric if (VariableSet.empty())
60406c3fb27SDimitry Andric return KernelSet;
605bdd1243dSDimitry Andric
606bdd1243dSDimitry Andric for (Function &Func : M.functions()) {
607bdd1243dSDimitry Andric if (Func.isDeclaration() || !isKernelLDS(&Func))
608bdd1243dSDimitry Andric continue;
609bdd1243dSDimitry Andric for (GlobalVariable *GV : LDSUsesInfo.indirect_access[&Func]) {
610bdd1243dSDimitry Andric if (VariableSet.contains(GV)) {
611bdd1243dSDimitry Andric KernelSet.insert(&Func);
612bdd1243dSDimitry Andric break;
613bdd1243dSDimitry Andric }
614bdd1243dSDimitry Andric }
615bdd1243dSDimitry Andric }
616bdd1243dSDimitry Andric
617bdd1243dSDimitry Andric return KernelSet;
618bdd1243dSDimitry Andric }
619bdd1243dSDimitry Andric
620bdd1243dSDimitry Andric static GlobalVariable *
chooseBestVariableForModuleStrategy(const DataLayout & DL,VariableFunctionMap & LDSVars)621bdd1243dSDimitry Andric chooseBestVariableForModuleStrategy(const DataLayout &DL,
622bdd1243dSDimitry Andric VariableFunctionMap &LDSVars) {
623bdd1243dSDimitry Andric // Find the global variable with the most indirect uses from kernels
624bdd1243dSDimitry Andric
625bdd1243dSDimitry Andric struct CandidateTy {
626bdd1243dSDimitry Andric GlobalVariable *GV = nullptr;
627bdd1243dSDimitry Andric size_t UserCount = 0;
628bdd1243dSDimitry Andric size_t Size = 0;
629bdd1243dSDimitry Andric
630bdd1243dSDimitry Andric CandidateTy() = default;
631bdd1243dSDimitry Andric
632bdd1243dSDimitry Andric CandidateTy(GlobalVariable *GV, uint64_t UserCount, uint64_t AllocSize)
633bdd1243dSDimitry Andric : GV(GV), UserCount(UserCount), Size(AllocSize) {}
634bdd1243dSDimitry Andric
635bdd1243dSDimitry Andric bool operator<(const CandidateTy &Other) const {
636bdd1243dSDimitry Andric // Fewer users makes module scope variable less attractive
637bdd1243dSDimitry Andric if (UserCount < Other.UserCount) {
638bdd1243dSDimitry Andric return true;
639bdd1243dSDimitry Andric }
640bdd1243dSDimitry Andric if (UserCount > Other.UserCount) {
641bdd1243dSDimitry Andric return false;
642bdd1243dSDimitry Andric }
643bdd1243dSDimitry Andric
644bdd1243dSDimitry Andric // Bigger makes module scope variable less attractive
645bdd1243dSDimitry Andric if (Size < Other.Size) {
646bdd1243dSDimitry Andric return false;
647bdd1243dSDimitry Andric }
648bdd1243dSDimitry Andric
649bdd1243dSDimitry Andric if (Size > Other.Size) {
650bdd1243dSDimitry Andric return true;
651bdd1243dSDimitry Andric }
652bdd1243dSDimitry Andric
653bdd1243dSDimitry Andric // Arbitrary but consistent
654bdd1243dSDimitry Andric return GV->getName() < Other.GV->getName();
655bdd1243dSDimitry Andric }
656bdd1243dSDimitry Andric };
657bdd1243dSDimitry Andric
658bdd1243dSDimitry Andric CandidateTy MostUsed;
659bdd1243dSDimitry Andric
660bdd1243dSDimitry Andric for (auto &K : LDSVars) {
661bdd1243dSDimitry Andric GlobalVariable *GV = K.first;
662bdd1243dSDimitry Andric if (K.second.size() <= 1) {
663bdd1243dSDimitry Andric // A variable reachable by only one kernel is best lowered with kernel
664bdd1243dSDimitry Andric // strategy
665bdd1243dSDimitry Andric continue;
666bdd1243dSDimitry Andric }
66706c3fb27SDimitry Andric CandidateTy Candidate(
66806c3fb27SDimitry Andric GV, K.second.size(),
669bdd1243dSDimitry Andric DL.getTypeAllocSize(GV->getValueType()).getFixedValue());
670bdd1243dSDimitry Andric if (MostUsed < Candidate)
671bdd1243dSDimitry Andric MostUsed = Candidate;
672bdd1243dSDimitry Andric }
673bdd1243dSDimitry Andric
674bdd1243dSDimitry Andric return MostUsed.GV;
675bdd1243dSDimitry Andric }
676bdd1243dSDimitry Andric
recordLDSAbsoluteAddress(Module * M,GlobalVariable * GV,uint32_t Address)67706c3fb27SDimitry Andric static void recordLDSAbsoluteAddress(Module *M, GlobalVariable *GV,
67806c3fb27SDimitry Andric uint32_t Address) {
67906c3fb27SDimitry Andric // Write the specified address into metadata where it can be retrieved by
68006c3fb27SDimitry Andric // the assembler. Format is a half open range, [Address Address+1)
68106c3fb27SDimitry Andric LLVMContext &Ctx = M->getContext();
68206c3fb27SDimitry Andric auto *IntTy =
68306c3fb27SDimitry Andric M->getDataLayout().getIntPtrType(Ctx, AMDGPUAS::LOCAL_ADDRESS);
68406c3fb27SDimitry Andric auto *MinC = ConstantAsMetadata::get(ConstantInt::get(IntTy, Address));
68506c3fb27SDimitry Andric auto *MaxC = ConstantAsMetadata::get(ConstantInt::get(IntTy, Address + 1));
68606c3fb27SDimitry Andric GV->setMetadata(LLVMContext::MD_absolute_symbol,
68706c3fb27SDimitry Andric MDNode::get(Ctx, {MinC, MaxC}));
68806c3fb27SDimitry Andric }
689972a253aSDimitry Andric
69006c3fb27SDimitry Andric DenseMap<Function *, Value *> tableKernelIndexCache;
getTableLookupKernelIndex(Module & M,Function * F)69106c3fb27SDimitry Andric Value *getTableLookupKernelIndex(Module &M, Function *F) {
69206c3fb27SDimitry Andric // Accesses from a function use the amdgcn_lds_kernel_id intrinsic which
69306c3fb27SDimitry Andric // lowers to a read from a live in register. Emit it once in the entry
69406c3fb27SDimitry Andric // block to spare deduplicating it later.
69506c3fb27SDimitry Andric auto [It, Inserted] = tableKernelIndexCache.try_emplace(F);
69606c3fb27SDimitry Andric if (Inserted) {
69706c3fb27SDimitry Andric Function *Decl =
69806c3fb27SDimitry Andric Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_lds_kernel_id, {});
699fe6060f1SDimitry Andric
70006c3fb27SDimitry Andric auto InsertAt = F->getEntryBlock().getFirstNonPHIOrDbgOrAlloca();
70106c3fb27SDimitry Andric IRBuilder<> Builder(&*InsertAt);
702972a253aSDimitry Andric
70306c3fb27SDimitry Andric It->second = Builder.CreateCall(Decl, {});
70406c3fb27SDimitry Andric }
705972a253aSDimitry Andric
70606c3fb27SDimitry Andric return It->second;
70706c3fb27SDimitry Andric }
70806c3fb27SDimitry Andric
assignLDSKernelIDToEachKernel(Module * M,DenseSet<Function * > const & KernelsThatAllocateTableLDS,DenseSet<Function * > const & KernelsThatIndirectlyAllocateDynamicLDS)70906c3fb27SDimitry Andric static std::vector<Function *> assignLDSKernelIDToEachKernel(
71006c3fb27SDimitry Andric Module *M, DenseSet<Function *> const &KernelsThatAllocateTableLDS,
71106c3fb27SDimitry Andric DenseSet<Function *> const &KernelsThatIndirectlyAllocateDynamicLDS) {
71206c3fb27SDimitry Andric // Associate kernels in the set with an arbirary but reproducible order and
71306c3fb27SDimitry Andric // annotate them with that order in metadata. This metadata is recognised by
71406c3fb27SDimitry Andric // the backend and lowered to a SGPR which can be read from using
71506c3fb27SDimitry Andric // amdgcn_lds_kernel_id.
71606c3fb27SDimitry Andric
71706c3fb27SDimitry Andric std::vector<Function *> OrderedKernels;
71806c3fb27SDimitry Andric if (!KernelsThatAllocateTableLDS.empty() ||
71906c3fb27SDimitry Andric !KernelsThatIndirectlyAllocateDynamicLDS.empty()) {
72006c3fb27SDimitry Andric
72106c3fb27SDimitry Andric for (Function &Func : M->functions()) {
72206c3fb27SDimitry Andric if (Func.isDeclaration())
72306c3fb27SDimitry Andric continue;
72406c3fb27SDimitry Andric if (!isKernelLDS(&Func))
72506c3fb27SDimitry Andric continue;
72606c3fb27SDimitry Andric
72706c3fb27SDimitry Andric if (KernelsThatAllocateTableLDS.contains(&Func) ||
72806c3fb27SDimitry Andric KernelsThatIndirectlyAllocateDynamicLDS.contains(&Func)) {
72906c3fb27SDimitry Andric assert(Func.hasName()); // else fatal error earlier
73006c3fb27SDimitry Andric OrderedKernels.push_back(&Func);
731bdd1243dSDimitry Andric }
732bdd1243dSDimitry Andric }
733972a253aSDimitry Andric
73406c3fb27SDimitry Andric // Put them in an arbitrary but reproducible order
73506c3fb27SDimitry Andric OrderedKernels = sortByName(std::move(OrderedKernels));
736972a253aSDimitry Andric
73706c3fb27SDimitry Andric // Annotate the kernels with their order in this vector
73806c3fb27SDimitry Andric LLVMContext &Ctx = M->getContext();
73906c3fb27SDimitry Andric IRBuilder<> Builder(Ctx);
74006c3fb27SDimitry Andric
74106c3fb27SDimitry Andric if (OrderedKernels.size() > UINT32_MAX) {
74206c3fb27SDimitry Andric // 32 bit keeps it in one SGPR. > 2**32 kernels won't fit on the GPU
74306c3fb27SDimitry Andric report_fatal_error("Unimplemented LDS lowering for > 2**32 kernels");
74406c3fb27SDimitry Andric }
74506c3fb27SDimitry Andric
74606c3fb27SDimitry Andric for (size_t i = 0; i < OrderedKernels.size(); i++) {
74706c3fb27SDimitry Andric Metadata *AttrMDArgs[1] = {
74806c3fb27SDimitry Andric ConstantAsMetadata::get(Builder.getInt32(i)),
74906c3fb27SDimitry Andric };
75006c3fb27SDimitry Andric OrderedKernels[i]->setMetadata("llvm.amdgcn.lds.kernel.id",
75106c3fb27SDimitry Andric MDNode::get(Ctx, AttrMDArgs));
75206c3fb27SDimitry Andric }
75306c3fb27SDimitry Andric }
75406c3fb27SDimitry Andric return OrderedKernels;
75506c3fb27SDimitry Andric }
75606c3fb27SDimitry Andric
partitionVariablesIntoIndirectStrategies(Module & M,LDSUsesInfoTy const & LDSUsesInfo,VariableFunctionMap & LDSToKernelsThatNeedToAccessItIndirectly,DenseSet<GlobalVariable * > & ModuleScopeVariables,DenseSet<GlobalVariable * > & TableLookupVariables,DenseSet<GlobalVariable * > & KernelAccessVariables,DenseSet<GlobalVariable * > & DynamicVariables)75706c3fb27SDimitry Andric static void partitionVariablesIntoIndirectStrategies(
75806c3fb27SDimitry Andric Module &M, LDSUsesInfoTy const &LDSUsesInfo,
75906c3fb27SDimitry Andric VariableFunctionMap &LDSToKernelsThatNeedToAccessItIndirectly,
76006c3fb27SDimitry Andric DenseSet<GlobalVariable *> &ModuleScopeVariables,
76106c3fb27SDimitry Andric DenseSet<GlobalVariable *> &TableLookupVariables,
76206c3fb27SDimitry Andric DenseSet<GlobalVariable *> &KernelAccessVariables,
76306c3fb27SDimitry Andric DenseSet<GlobalVariable *> &DynamicVariables) {
76406c3fb27SDimitry Andric
765bdd1243dSDimitry Andric GlobalVariable *HybridModuleRoot =
766bdd1243dSDimitry Andric LoweringKindLoc != LoweringKind::hybrid
767bdd1243dSDimitry Andric ? nullptr
768bdd1243dSDimitry Andric : chooseBestVariableForModuleStrategy(
76906c3fb27SDimitry Andric M.getDataLayout(), LDSToKernelsThatNeedToAccessItIndirectly);
770972a253aSDimitry Andric
771bdd1243dSDimitry Andric DenseSet<Function *> const EmptySet;
772bdd1243dSDimitry Andric DenseSet<Function *> const &HybridModuleRootKernels =
773bdd1243dSDimitry Andric HybridModuleRoot
774bdd1243dSDimitry Andric ? LDSToKernelsThatNeedToAccessItIndirectly[HybridModuleRoot]
775bdd1243dSDimitry Andric : EmptySet;
776bdd1243dSDimitry Andric
777bdd1243dSDimitry Andric for (auto &K : LDSToKernelsThatNeedToAccessItIndirectly) {
778bdd1243dSDimitry Andric // Each iteration of this loop assigns exactly one global variable to
779bdd1243dSDimitry Andric // exactly one of the implementation strategies.
780bdd1243dSDimitry Andric
781bdd1243dSDimitry Andric GlobalVariable *GV = K.first;
782bdd1243dSDimitry Andric assert(AMDGPU::isLDSVariableToLower(*GV));
783bdd1243dSDimitry Andric assert(K.second.size() != 0);
784bdd1243dSDimitry Andric
78506c3fb27SDimitry Andric if (AMDGPU::isDynamicLDS(*GV)) {
78606c3fb27SDimitry Andric DynamicVariables.insert(GV);
78706c3fb27SDimitry Andric continue;
78806c3fb27SDimitry Andric }
78906c3fb27SDimitry Andric
790bdd1243dSDimitry Andric switch (LoweringKindLoc) {
791bdd1243dSDimitry Andric case LoweringKind::module:
792bdd1243dSDimitry Andric ModuleScopeVariables.insert(GV);
793bdd1243dSDimitry Andric break;
794bdd1243dSDimitry Andric
795bdd1243dSDimitry Andric case LoweringKind::table:
796bdd1243dSDimitry Andric TableLookupVariables.insert(GV);
797bdd1243dSDimitry Andric break;
798bdd1243dSDimitry Andric
799bdd1243dSDimitry Andric case LoweringKind::kernel:
800bdd1243dSDimitry Andric if (K.second.size() == 1) {
801bdd1243dSDimitry Andric KernelAccessVariables.insert(GV);
802972a253aSDimitry Andric } else {
803bdd1243dSDimitry Andric report_fatal_error(
804bdd1243dSDimitry Andric "cannot lower LDS '" + GV->getName() +
805bdd1243dSDimitry Andric "' to kernel access as it is reachable from multiple kernels");
806bdd1243dSDimitry Andric }
807bdd1243dSDimitry Andric break;
808bdd1243dSDimitry Andric
809bdd1243dSDimitry Andric case LoweringKind::hybrid: {
810bdd1243dSDimitry Andric if (GV == HybridModuleRoot) {
811bdd1243dSDimitry Andric assert(K.second.size() != 1);
812bdd1243dSDimitry Andric ModuleScopeVariables.insert(GV);
813bdd1243dSDimitry Andric } else if (K.second.size() == 1) {
814bdd1243dSDimitry Andric KernelAccessVariables.insert(GV);
815bdd1243dSDimitry Andric } else if (set_is_subset(K.second, HybridModuleRootKernels)) {
816bdd1243dSDimitry Andric ModuleScopeVariables.insert(GV);
817bdd1243dSDimitry Andric } else {
818bdd1243dSDimitry Andric TableLookupVariables.insert(GV);
819bdd1243dSDimitry Andric }
820bdd1243dSDimitry Andric break;
821bdd1243dSDimitry Andric }
822bdd1243dSDimitry Andric }
823bdd1243dSDimitry Andric }
824bdd1243dSDimitry Andric
82506c3fb27SDimitry Andric // All LDS variables accessed indirectly have now been partitioned into
82606c3fb27SDimitry Andric // the distinct lowering strategies.
827bdd1243dSDimitry Andric assert(ModuleScopeVariables.size() + TableLookupVariables.size() +
82806c3fb27SDimitry Andric KernelAccessVariables.size() + DynamicVariables.size() ==
829bdd1243dSDimitry Andric LDSToKernelsThatNeedToAccessItIndirectly.size());
83006c3fb27SDimitry Andric }
831bdd1243dSDimitry Andric
lowerModuleScopeStructVariables(Module & M,DenseSet<GlobalVariable * > const & ModuleScopeVariables,DenseSet<Function * > const & KernelsThatAllocateModuleLDS)83206c3fb27SDimitry Andric static GlobalVariable *lowerModuleScopeStructVariables(
83306c3fb27SDimitry Andric Module &M, DenseSet<GlobalVariable *> const &ModuleScopeVariables,
83406c3fb27SDimitry Andric DenseSet<Function *> const &KernelsThatAllocateModuleLDS) {
83506c3fb27SDimitry Andric // Create a struct to hold the ModuleScopeVariables
83606c3fb27SDimitry Andric // Replace all uses of those variables from non-kernel functions with the
83706c3fb27SDimitry Andric // new struct instance Replace only the uses from kernel functions that will
83806c3fb27SDimitry Andric // allocate this instance. That is a space optimisation - kernels that use a
83906c3fb27SDimitry Andric // subset of the module scope struct and do not need to allocate it for
84006c3fb27SDimitry Andric // indirect calls will only allocate the subset they use (they do so as part
84106c3fb27SDimitry Andric // of the per-kernel lowering).
84206c3fb27SDimitry Andric if (ModuleScopeVariables.empty()) {
84306c3fb27SDimitry Andric return nullptr;
84406c3fb27SDimitry Andric }
845bdd1243dSDimitry Andric
84606c3fb27SDimitry Andric LLVMContext &Ctx = M.getContext();
84706c3fb27SDimitry Andric
848bdd1243dSDimitry Andric LDSVariableReplacement ModuleScopeReplacement =
849bdd1243dSDimitry Andric createLDSVariableReplacement(M, "llvm.amdgcn.module.lds",
850bdd1243dSDimitry Andric ModuleScopeVariables);
851bdd1243dSDimitry Andric
85206c3fb27SDimitry Andric appendToCompilerUsed(M, {static_cast<GlobalValue *>(
853bdd1243dSDimitry Andric ConstantExpr::getPointerBitCastOrAddrSpaceCast(
854bdd1243dSDimitry Andric cast<Constant>(ModuleScopeReplacement.SGV),
8555f757f3fSDimitry Andric PointerType::getUnqual(Ctx)))});
856bdd1243dSDimitry Andric
85706c3fb27SDimitry Andric // module.lds will be allocated at zero in any kernel that allocates it
85806c3fb27SDimitry Andric recordLDSAbsoluteAddress(&M, ModuleScopeReplacement.SGV, 0);
85906c3fb27SDimitry Andric
860bdd1243dSDimitry Andric // historic
861bdd1243dSDimitry Andric removeLocalVarsFromUsedLists(M, ModuleScopeVariables);
862bdd1243dSDimitry Andric
863bdd1243dSDimitry Andric // Replace all uses of module scope variable from non-kernel functions
864bdd1243dSDimitry Andric replaceLDSVariablesWithStruct(
865bdd1243dSDimitry Andric M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) {
866bdd1243dSDimitry Andric Instruction *I = dyn_cast<Instruction>(U.getUser());
867bdd1243dSDimitry Andric if (!I) {
868bdd1243dSDimitry Andric return false;
869bdd1243dSDimitry Andric }
870bdd1243dSDimitry Andric Function *F = I->getFunction();
871bdd1243dSDimitry Andric return !isKernelLDS(F);
872bdd1243dSDimitry Andric });
873bdd1243dSDimitry Andric
874bdd1243dSDimitry Andric // Replace uses of module scope variable from kernel functions that
875bdd1243dSDimitry Andric // allocate the module scope variable, otherwise leave them unchanged
876bdd1243dSDimitry Andric // Record on each kernel whether the module scope global is used by it
877bdd1243dSDimitry Andric
878bdd1243dSDimitry Andric for (Function &Func : M.functions()) {
879bdd1243dSDimitry Andric if (Func.isDeclaration() || !isKernelLDS(&Func))
880bdd1243dSDimitry Andric continue;
881bdd1243dSDimitry Andric
882bdd1243dSDimitry Andric if (KernelsThatAllocateModuleLDS.contains(&Func)) {
883bdd1243dSDimitry Andric replaceLDSVariablesWithStruct(
884bdd1243dSDimitry Andric M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) {
885bdd1243dSDimitry Andric Instruction *I = dyn_cast<Instruction>(U.getUser());
886bdd1243dSDimitry Andric if (!I) {
887bdd1243dSDimitry Andric return false;
888bdd1243dSDimitry Andric }
889bdd1243dSDimitry Andric Function *F = I->getFunction();
890bdd1243dSDimitry Andric return F == &Func;
891bdd1243dSDimitry Andric });
892bdd1243dSDimitry Andric
89306c3fb27SDimitry Andric markUsedByKernel(&Func, ModuleScopeReplacement.SGV);
894972a253aSDimitry Andric }
895972a253aSDimitry Andric }
896972a253aSDimitry Andric
89706c3fb27SDimitry Andric return ModuleScopeReplacement.SGV;
89806c3fb27SDimitry Andric }
89906c3fb27SDimitry Andric
90006c3fb27SDimitry Andric static DenseMap<Function *, LDSVariableReplacement>
lowerKernelScopeStructVariables(Module & M,LDSUsesInfoTy & LDSUsesInfo,DenseSet<GlobalVariable * > const & ModuleScopeVariables,DenseSet<Function * > const & KernelsThatAllocateModuleLDS,GlobalVariable * MaybeModuleScopeStruct)90106c3fb27SDimitry Andric lowerKernelScopeStructVariables(
90206c3fb27SDimitry Andric Module &M, LDSUsesInfoTy &LDSUsesInfo,
90306c3fb27SDimitry Andric DenseSet<GlobalVariable *> const &ModuleScopeVariables,
90406c3fb27SDimitry Andric DenseSet<Function *> const &KernelsThatAllocateModuleLDS,
90506c3fb27SDimitry Andric GlobalVariable *MaybeModuleScopeStruct) {
90606c3fb27SDimitry Andric
90706c3fb27SDimitry Andric // Create a struct for each kernel for the non-module-scope variables.
90806c3fb27SDimitry Andric
909bdd1243dSDimitry Andric DenseMap<Function *, LDSVariableReplacement> KernelToReplacement;
910bdd1243dSDimitry Andric for (Function &Func : M.functions()) {
911bdd1243dSDimitry Andric if (Func.isDeclaration() || !isKernelLDS(&Func))
912349cc55cSDimitry Andric continue;
913349cc55cSDimitry Andric
914bdd1243dSDimitry Andric DenseSet<GlobalVariable *> KernelUsedVariables;
91506c3fb27SDimitry Andric // Allocating variables that are used directly in this struct to get
91606c3fb27SDimitry Andric // alignment aware allocation and predictable frame size.
917bdd1243dSDimitry Andric for (auto &v : LDSUsesInfo.direct_access[&Func]) {
91806c3fb27SDimitry Andric if (!AMDGPU::isDynamicLDS(*v)) {
919bdd1243dSDimitry Andric KernelUsedVariables.insert(v);
920bdd1243dSDimitry Andric }
92106c3fb27SDimitry Andric }
92206c3fb27SDimitry Andric
92306c3fb27SDimitry Andric // Allocating variables that are accessed indirectly so that a lookup of
92406c3fb27SDimitry Andric // this struct instance can find them from nested functions.
925bdd1243dSDimitry Andric for (auto &v : LDSUsesInfo.indirect_access[&Func]) {
92606c3fb27SDimitry Andric if (!AMDGPU::isDynamicLDS(*v)) {
927bdd1243dSDimitry Andric KernelUsedVariables.insert(v);
928bdd1243dSDimitry Andric }
92906c3fb27SDimitry Andric }
930bdd1243dSDimitry Andric
931bdd1243dSDimitry Andric // Variables allocated in module lds must all resolve to that struct,
932bdd1243dSDimitry Andric // not to the per-kernel instance.
933bdd1243dSDimitry Andric if (KernelsThatAllocateModuleLDS.contains(&Func)) {
934bdd1243dSDimitry Andric for (GlobalVariable *v : ModuleScopeVariables) {
935bdd1243dSDimitry Andric KernelUsedVariables.erase(v);
936bdd1243dSDimitry Andric }
937bdd1243dSDimitry Andric }
938bdd1243dSDimitry Andric
939bdd1243dSDimitry Andric if (KernelUsedVariables.empty()) {
94006c3fb27SDimitry Andric // Either used no LDS, or the LDS it used was all in the module struct
94106c3fb27SDimitry Andric // or dynamically sized
942fe6060f1SDimitry Andric continue;
943972a253aSDimitry Andric }
944972a253aSDimitry Andric
945bdd1243dSDimitry Andric // The association between kernel function and LDS struct is done by
946bdd1243dSDimitry Andric // symbol name, which only works if the function in question has a
947bdd1243dSDimitry Andric // name This is not expected to be a problem in practice as kernels
948bdd1243dSDimitry Andric // are called by name making anonymous ones (which are named by the
949bdd1243dSDimitry Andric // backend) difficult to use. This does mean that llvm test cases need
950bdd1243dSDimitry Andric // to name the kernels.
951bdd1243dSDimitry Andric if (!Func.hasName()) {
952bdd1243dSDimitry Andric report_fatal_error("Anonymous kernels cannot use LDS variables");
953bdd1243dSDimitry Andric }
954bdd1243dSDimitry Andric
955972a253aSDimitry Andric std::string VarName =
956bdd1243dSDimitry Andric (Twine("llvm.amdgcn.kernel.") + Func.getName() + ".lds").str();
957bdd1243dSDimitry Andric
958bdd1243dSDimitry Andric auto Replacement =
959972a253aSDimitry Andric createLDSVariableReplacement(M, VarName, KernelUsedVariables);
960972a253aSDimitry Andric
96106c3fb27SDimitry Andric // If any indirect uses, create a direct use to ensure allocation
96206c3fb27SDimitry Andric // TODO: Simpler to unconditionally mark used but that regresses
96306c3fb27SDimitry Andric // codegen in test/CodeGen/AMDGPU/noclobber-barrier.ll
96406c3fb27SDimitry Andric auto Accesses = LDSUsesInfo.indirect_access.find(&Func);
96506c3fb27SDimitry Andric if ((Accesses != LDSUsesInfo.indirect_access.end()) &&
96606c3fb27SDimitry Andric !Accesses->second.empty())
96706c3fb27SDimitry Andric markUsedByKernel(&Func, Replacement.SGV);
96806c3fb27SDimitry Andric
969bdd1243dSDimitry Andric // remove preserves existing codegen
970bdd1243dSDimitry Andric removeLocalVarsFromUsedLists(M, KernelUsedVariables);
971bdd1243dSDimitry Andric KernelToReplacement[&Func] = Replacement;
972bdd1243dSDimitry Andric
973bdd1243dSDimitry Andric // Rewrite uses within kernel to the new struct
974972a253aSDimitry Andric replaceLDSVariablesWithStruct(
975bdd1243dSDimitry Andric M, KernelUsedVariables, Replacement, [&Func](Use &U) {
976972a253aSDimitry Andric Instruction *I = dyn_cast<Instruction>(U.getUser());
977bdd1243dSDimitry Andric return I && I->getFunction() == &Func;
978972a253aSDimitry Andric });
979972a253aSDimitry Andric }
98006c3fb27SDimitry Andric return KernelToReplacement;
98106c3fb27SDimitry Andric }
98206c3fb27SDimitry Andric
98306c3fb27SDimitry Andric static GlobalVariable *
buildRepresentativeDynamicLDSInstance(Module & M,LDSUsesInfoTy & LDSUsesInfo,Function * func)98406c3fb27SDimitry Andric buildRepresentativeDynamicLDSInstance(Module &M, LDSUsesInfoTy &LDSUsesInfo,
98506c3fb27SDimitry Andric Function *func) {
98606c3fb27SDimitry Andric // Create a dynamic lds variable with a name associated with the passed
98706c3fb27SDimitry Andric // function that has the maximum alignment of any dynamic lds variable
98806c3fb27SDimitry Andric // reachable from this kernel. Dynamic LDS is allocated after the static LDS
98906c3fb27SDimitry Andric // allocation, possibly after alignment padding. The representative variable
99006c3fb27SDimitry Andric // created here has the maximum alignment of any other dynamic variable
99106c3fb27SDimitry Andric // reachable by that kernel. All dynamic LDS variables are allocated at the
99206c3fb27SDimitry Andric // same address in each kernel in order to provide the documented aliasing
99306c3fb27SDimitry Andric // semantics. Setting the alignment here allows this IR pass to accurately
99406c3fb27SDimitry Andric // predict the exact constant at which it will be allocated.
99506c3fb27SDimitry Andric
99606c3fb27SDimitry Andric assert(isKernelLDS(func));
99706c3fb27SDimitry Andric
99806c3fb27SDimitry Andric LLVMContext &Ctx = M.getContext();
99906c3fb27SDimitry Andric const DataLayout &DL = M.getDataLayout();
100006c3fb27SDimitry Andric Align MaxDynamicAlignment(1);
100106c3fb27SDimitry Andric
100206c3fb27SDimitry Andric auto UpdateMaxAlignment = [&MaxDynamicAlignment, &DL](GlobalVariable *GV) {
100306c3fb27SDimitry Andric if (AMDGPU::isDynamicLDS(*GV)) {
100406c3fb27SDimitry Andric MaxDynamicAlignment =
100506c3fb27SDimitry Andric std::max(MaxDynamicAlignment, AMDGPU::getAlign(DL, GV));
100606c3fb27SDimitry Andric }
100706c3fb27SDimitry Andric };
100806c3fb27SDimitry Andric
100906c3fb27SDimitry Andric for (GlobalVariable *GV : LDSUsesInfo.indirect_access[func]) {
101006c3fb27SDimitry Andric UpdateMaxAlignment(GV);
101106c3fb27SDimitry Andric }
101206c3fb27SDimitry Andric
101306c3fb27SDimitry Andric for (GlobalVariable *GV : LDSUsesInfo.direct_access[func]) {
101406c3fb27SDimitry Andric UpdateMaxAlignment(GV);
101506c3fb27SDimitry Andric }
101606c3fb27SDimitry Andric
101706c3fb27SDimitry Andric assert(func->hasName()); // Checked by caller
101806c3fb27SDimitry Andric auto emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);
101906c3fb27SDimitry Andric GlobalVariable *N = new GlobalVariable(
102006c3fb27SDimitry Andric M, emptyCharArray, false, GlobalValue::ExternalLinkage, nullptr,
102106c3fb27SDimitry Andric Twine("llvm.amdgcn." + func->getName() + ".dynlds"), nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS,
102206c3fb27SDimitry Andric false);
102306c3fb27SDimitry Andric N->setAlignment(MaxDynamicAlignment);
102406c3fb27SDimitry Andric
102506c3fb27SDimitry Andric assert(AMDGPU::isDynamicLDS(*N));
102606c3fb27SDimitry Andric return N;
102706c3fb27SDimitry Andric }
102806c3fb27SDimitry Andric
1029*297eecfbSDimitry Andric /// Strip "amdgpu-no-lds-kernel-id" from any functions where we may have
1030*297eecfbSDimitry Andric /// introduced its use. If AMDGPUAttributor ran prior to the pass, we inferred
1031*297eecfbSDimitry Andric /// the lack of llvm.amdgcn.lds.kernel.id calls.
removeNoLdsKernelIdFromReachable(CallGraph & CG,Function * KernelRoot)1032*297eecfbSDimitry Andric void removeNoLdsKernelIdFromReachable(CallGraph &CG, Function *KernelRoot) {
1033*297eecfbSDimitry Andric KernelRoot->removeFnAttr("amdgpu-no-lds-kernel-id");
1034*297eecfbSDimitry Andric
1035*297eecfbSDimitry Andric SmallVector<Function *> Tmp({CG[KernelRoot]->getFunction()});
1036*297eecfbSDimitry Andric if (!Tmp.back())
1037*297eecfbSDimitry Andric return;
1038*297eecfbSDimitry Andric
1039*297eecfbSDimitry Andric SmallPtrSet<Function *, 8> Visited;
1040*297eecfbSDimitry Andric bool SeenUnknownCall = false;
1041*297eecfbSDimitry Andric
1042*297eecfbSDimitry Andric do {
1043*297eecfbSDimitry Andric Function *F = Tmp.pop_back_val();
1044*297eecfbSDimitry Andric
1045*297eecfbSDimitry Andric for (auto &N : *CG[F]) {
1046*297eecfbSDimitry Andric if (!N.second)
1047*297eecfbSDimitry Andric continue;
1048*297eecfbSDimitry Andric
1049*297eecfbSDimitry Andric Function *Callee = N.second->getFunction();
1050*297eecfbSDimitry Andric if (!Callee) {
1051*297eecfbSDimitry Andric if (!SeenUnknownCall) {
1052*297eecfbSDimitry Andric SeenUnknownCall = true;
1053*297eecfbSDimitry Andric
1054*297eecfbSDimitry Andric // If we see any indirect calls, assume nothing about potential
1055*297eecfbSDimitry Andric // targets.
1056*297eecfbSDimitry Andric // TODO: This could be refined to possible LDS global users.
1057*297eecfbSDimitry Andric for (auto &N : *CG.getExternalCallingNode()) {
1058*297eecfbSDimitry Andric Function *PotentialCallee = N.second->getFunction();
1059*297eecfbSDimitry Andric if (!isKernelLDS(PotentialCallee))
1060*297eecfbSDimitry Andric PotentialCallee->removeFnAttr("amdgpu-no-lds-kernel-id");
1061*297eecfbSDimitry Andric }
1062*297eecfbSDimitry Andric
1063*297eecfbSDimitry Andric continue;
1064*297eecfbSDimitry Andric }
1065*297eecfbSDimitry Andric }
1066*297eecfbSDimitry Andric
1067*297eecfbSDimitry Andric Callee->removeFnAttr("amdgpu-no-lds-kernel-id");
1068*297eecfbSDimitry Andric if (Visited.insert(Callee).second)
1069*297eecfbSDimitry Andric Tmp.push_back(Callee);
1070*297eecfbSDimitry Andric }
1071*297eecfbSDimitry Andric } while (!Tmp.empty());
1072*297eecfbSDimitry Andric }
1073*297eecfbSDimitry Andric
lowerDynamicLDSVariables(Module & M,LDSUsesInfoTy & LDSUsesInfo,DenseSet<Function * > const & KernelsThatIndirectlyAllocateDynamicLDS,DenseSet<GlobalVariable * > const & DynamicVariables,std::vector<Function * > const & OrderedKernels)107406c3fb27SDimitry Andric DenseMap<Function *, GlobalVariable *> lowerDynamicLDSVariables(
107506c3fb27SDimitry Andric Module &M, LDSUsesInfoTy &LDSUsesInfo,
107606c3fb27SDimitry Andric DenseSet<Function *> const &KernelsThatIndirectlyAllocateDynamicLDS,
107706c3fb27SDimitry Andric DenseSet<GlobalVariable *> const &DynamicVariables,
107806c3fb27SDimitry Andric std::vector<Function *> const &OrderedKernels) {
107906c3fb27SDimitry Andric DenseMap<Function *, GlobalVariable *> KernelToCreatedDynamicLDS;
108006c3fb27SDimitry Andric if (!KernelsThatIndirectlyAllocateDynamicLDS.empty()) {
108106c3fb27SDimitry Andric LLVMContext &Ctx = M.getContext();
108206c3fb27SDimitry Andric IRBuilder<> Builder(Ctx);
108306c3fb27SDimitry Andric Type *I32 = Type::getInt32Ty(Ctx);
108406c3fb27SDimitry Andric
108506c3fb27SDimitry Andric std::vector<Constant *> newDynamicLDS;
108606c3fb27SDimitry Andric
108706c3fb27SDimitry Andric // Table is built in the same order as OrderedKernels
108806c3fb27SDimitry Andric for (auto &func : OrderedKernels) {
108906c3fb27SDimitry Andric
109006c3fb27SDimitry Andric if (KernelsThatIndirectlyAllocateDynamicLDS.contains(func)) {
109106c3fb27SDimitry Andric assert(isKernelLDS(func));
109206c3fb27SDimitry Andric if (!func->hasName()) {
109306c3fb27SDimitry Andric report_fatal_error("Anonymous kernels cannot use LDS variables");
109406c3fb27SDimitry Andric }
109506c3fb27SDimitry Andric
109606c3fb27SDimitry Andric GlobalVariable *N =
109706c3fb27SDimitry Andric buildRepresentativeDynamicLDSInstance(M, LDSUsesInfo, func);
109806c3fb27SDimitry Andric
109906c3fb27SDimitry Andric KernelToCreatedDynamicLDS[func] = N;
110006c3fb27SDimitry Andric
110106c3fb27SDimitry Andric markUsedByKernel(func, N);
110206c3fb27SDimitry Andric
110306c3fb27SDimitry Andric auto emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);
110406c3fb27SDimitry Andric auto GEP = ConstantExpr::getGetElementPtr(
110506c3fb27SDimitry Andric emptyCharArray, N, ConstantInt::get(I32, 0), true);
110606c3fb27SDimitry Andric newDynamicLDS.push_back(ConstantExpr::getPtrToInt(GEP, I32));
110706c3fb27SDimitry Andric } else {
110806c3fb27SDimitry Andric newDynamicLDS.push_back(PoisonValue::get(I32));
110906c3fb27SDimitry Andric }
111006c3fb27SDimitry Andric }
111106c3fb27SDimitry Andric assert(OrderedKernels.size() == newDynamicLDS.size());
111206c3fb27SDimitry Andric
111306c3fb27SDimitry Andric ArrayType *t = ArrayType::get(I32, newDynamicLDS.size());
111406c3fb27SDimitry Andric Constant *init = ConstantArray::get(t, newDynamicLDS);
111506c3fb27SDimitry Andric GlobalVariable *table = new GlobalVariable(
111606c3fb27SDimitry Andric M, t, true, GlobalValue::InternalLinkage, init,
111706c3fb27SDimitry Andric "llvm.amdgcn.dynlds.offset.table", nullptr,
111806c3fb27SDimitry Andric GlobalValue::NotThreadLocal, AMDGPUAS::CONSTANT_ADDRESS);
111906c3fb27SDimitry Andric
112006c3fb27SDimitry Andric for (GlobalVariable *GV : DynamicVariables) {
112106c3fb27SDimitry Andric for (Use &U : make_early_inc_range(GV->uses())) {
112206c3fb27SDimitry Andric auto *I = dyn_cast<Instruction>(U.getUser());
112306c3fb27SDimitry Andric if (!I)
112406c3fb27SDimitry Andric continue;
112506c3fb27SDimitry Andric if (isKernelLDS(I->getFunction()))
112606c3fb27SDimitry Andric continue;
112706c3fb27SDimitry Andric
112806c3fb27SDimitry Andric replaceUseWithTableLookup(M, Builder, table, GV, U, nullptr);
112906c3fb27SDimitry Andric }
113006c3fb27SDimitry Andric }
113106c3fb27SDimitry Andric }
113206c3fb27SDimitry Andric return KernelToCreatedDynamicLDS;
113306c3fb27SDimitry Andric }
113406c3fb27SDimitry Andric
runOnModule(Module & M)11355f757f3fSDimitry Andric bool runOnModule(Module &M) {
113606c3fb27SDimitry Andric CallGraph CG = CallGraph(M);
113706c3fb27SDimitry Andric bool Changed = superAlignLDSGlobals(M);
113806c3fb27SDimitry Andric
113906c3fb27SDimitry Andric Changed |= eliminateConstantExprUsesOfLDSFromAllInstructions(M);
114006c3fb27SDimitry Andric
114106c3fb27SDimitry Andric Changed = true; // todo: narrow this down
114206c3fb27SDimitry Andric
114306c3fb27SDimitry Andric // For each kernel, what variables does it access directly or through
114406c3fb27SDimitry Andric // callees
114506c3fb27SDimitry Andric LDSUsesInfoTy LDSUsesInfo = getTransitiveUsesOfLDS(CG, M);
114606c3fb27SDimitry Andric
114706c3fb27SDimitry Andric // For each variable accessed through callees, which kernels access it
114806c3fb27SDimitry Andric VariableFunctionMap LDSToKernelsThatNeedToAccessItIndirectly;
114906c3fb27SDimitry Andric for (auto &K : LDSUsesInfo.indirect_access) {
115006c3fb27SDimitry Andric Function *F = K.first;
115106c3fb27SDimitry Andric assert(isKernelLDS(F));
115206c3fb27SDimitry Andric for (GlobalVariable *GV : K.second) {
115306c3fb27SDimitry Andric LDSToKernelsThatNeedToAccessItIndirectly[GV].insert(F);
115406c3fb27SDimitry Andric }
115506c3fb27SDimitry Andric }
115606c3fb27SDimitry Andric
115706c3fb27SDimitry Andric // Partition variables accessed indirectly into the different strategies
115806c3fb27SDimitry Andric DenseSet<GlobalVariable *> ModuleScopeVariables;
115906c3fb27SDimitry Andric DenseSet<GlobalVariable *> TableLookupVariables;
116006c3fb27SDimitry Andric DenseSet<GlobalVariable *> KernelAccessVariables;
116106c3fb27SDimitry Andric DenseSet<GlobalVariable *> DynamicVariables;
116206c3fb27SDimitry Andric partitionVariablesIntoIndirectStrategies(
116306c3fb27SDimitry Andric M, LDSUsesInfo, LDSToKernelsThatNeedToAccessItIndirectly,
116406c3fb27SDimitry Andric ModuleScopeVariables, TableLookupVariables, KernelAccessVariables,
116506c3fb27SDimitry Andric DynamicVariables);
116606c3fb27SDimitry Andric
116706c3fb27SDimitry Andric // If the kernel accesses a variable that is going to be stored in the
116806c3fb27SDimitry Andric // module instance through a call then that kernel needs to allocate the
116906c3fb27SDimitry Andric // module instance
117006c3fb27SDimitry Andric const DenseSet<Function *> KernelsThatAllocateModuleLDS =
117106c3fb27SDimitry Andric kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo,
117206c3fb27SDimitry Andric ModuleScopeVariables);
117306c3fb27SDimitry Andric const DenseSet<Function *> KernelsThatAllocateTableLDS =
117406c3fb27SDimitry Andric kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo,
117506c3fb27SDimitry Andric TableLookupVariables);
117606c3fb27SDimitry Andric
117706c3fb27SDimitry Andric const DenseSet<Function *> KernelsThatIndirectlyAllocateDynamicLDS =
117806c3fb27SDimitry Andric kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo,
117906c3fb27SDimitry Andric DynamicVariables);
118006c3fb27SDimitry Andric
118106c3fb27SDimitry Andric GlobalVariable *MaybeModuleScopeStruct = lowerModuleScopeStructVariables(
118206c3fb27SDimitry Andric M, ModuleScopeVariables, KernelsThatAllocateModuleLDS);
118306c3fb27SDimitry Andric
118406c3fb27SDimitry Andric DenseMap<Function *, LDSVariableReplacement> KernelToReplacement =
118506c3fb27SDimitry Andric lowerKernelScopeStructVariables(M, LDSUsesInfo, ModuleScopeVariables,
118606c3fb27SDimitry Andric KernelsThatAllocateModuleLDS,
118706c3fb27SDimitry Andric MaybeModuleScopeStruct);
1188bdd1243dSDimitry Andric
1189bdd1243dSDimitry Andric // Lower zero cost accesses to the kernel instances just created
1190bdd1243dSDimitry Andric for (auto &GV : KernelAccessVariables) {
1191bdd1243dSDimitry Andric auto &funcs = LDSToKernelsThatNeedToAccessItIndirectly[GV];
1192bdd1243dSDimitry Andric assert(funcs.size() == 1); // Only one kernel can access it
1193bdd1243dSDimitry Andric LDSVariableReplacement Replacement =
1194bdd1243dSDimitry Andric KernelToReplacement[*(funcs.begin())];
1195bdd1243dSDimitry Andric
1196bdd1243dSDimitry Andric DenseSet<GlobalVariable *> Vec;
1197bdd1243dSDimitry Andric Vec.insert(GV);
1198bdd1243dSDimitry Andric
1199bdd1243dSDimitry Andric replaceLDSVariablesWithStruct(M, Vec, Replacement, [](Use &U) {
1200bdd1243dSDimitry Andric return isa<Instruction>(U.getUser());
1201bdd1243dSDimitry Andric });
1202bdd1243dSDimitry Andric }
1203bdd1243dSDimitry Andric
120406c3fb27SDimitry Andric // The ith element of this vector is kernel id i
120506c3fb27SDimitry Andric std::vector<Function *> OrderedKernels =
120606c3fb27SDimitry Andric assignLDSKernelIDToEachKernel(&M, KernelsThatAllocateTableLDS,
120706c3fb27SDimitry Andric KernelsThatIndirectlyAllocateDynamicLDS);
120806c3fb27SDimitry Andric
1209bdd1243dSDimitry Andric if (!KernelsThatAllocateTableLDS.empty()) {
1210bdd1243dSDimitry Andric LLVMContext &Ctx = M.getContext();
1211bdd1243dSDimitry Andric IRBuilder<> Builder(Ctx);
1212bdd1243dSDimitry Andric
1213bdd1243dSDimitry Andric // The order must be consistent between lookup table and accesses to
1214bdd1243dSDimitry Andric // lookup table
121506c3fb27SDimitry Andric auto TableLookupVariablesOrdered =
121606c3fb27SDimitry Andric sortByName(std::vector<GlobalVariable *>(TableLookupVariables.begin(),
121706c3fb27SDimitry Andric TableLookupVariables.end()));
1218bdd1243dSDimitry Andric
1219bdd1243dSDimitry Andric GlobalVariable *LookupTable = buildLookupTable(
1220bdd1243dSDimitry Andric M, TableLookupVariablesOrdered, OrderedKernels, KernelToReplacement);
1221bdd1243dSDimitry Andric replaceUsesInInstructionsWithTableLookup(M, TableLookupVariablesOrdered,
1222bdd1243dSDimitry Andric LookupTable);
1223*297eecfbSDimitry Andric
1224*297eecfbSDimitry Andric // Strip amdgpu-no-lds-kernel-id from all functions reachable from the
1225*297eecfbSDimitry Andric // kernel. We may have inferred this wasn't used prior to the pass.
1226*297eecfbSDimitry Andric //
1227*297eecfbSDimitry Andric // TODO: We could filter out subgraphs that do not access LDS globals.
1228*297eecfbSDimitry Andric for (Function *F : KernelsThatAllocateTableLDS)
1229*297eecfbSDimitry Andric removeNoLdsKernelIdFromReachable(CG, F);
1230bdd1243dSDimitry Andric }
1231bdd1243dSDimitry Andric
123206c3fb27SDimitry Andric DenseMap<Function *, GlobalVariable *> KernelToCreatedDynamicLDS =
123306c3fb27SDimitry Andric lowerDynamicLDSVariables(M, LDSUsesInfo,
123406c3fb27SDimitry Andric KernelsThatIndirectlyAllocateDynamicLDS,
123506c3fb27SDimitry Andric DynamicVariables, OrderedKernels);
123606c3fb27SDimitry Andric
123706c3fb27SDimitry Andric // All kernel frames have been allocated. Calculate and record the
123806c3fb27SDimitry Andric // addresses.
123906c3fb27SDimitry Andric {
124006c3fb27SDimitry Andric const DataLayout &DL = M.getDataLayout();
124106c3fb27SDimitry Andric
124206c3fb27SDimitry Andric for (Function &Func : M.functions()) {
124306c3fb27SDimitry Andric if (Func.isDeclaration() || !isKernelLDS(&Func))
124406c3fb27SDimitry Andric continue;
124506c3fb27SDimitry Andric
124606c3fb27SDimitry Andric // All three of these are optional. The first variable is allocated at
124706c3fb27SDimitry Andric // zero. They are allocated by AMDGPUMachineFunction as one block.
124806c3fb27SDimitry Andric // Layout:
124906c3fb27SDimitry Andric //{
125006c3fb27SDimitry Andric // module.lds
125106c3fb27SDimitry Andric // alignment padding
125206c3fb27SDimitry Andric // kernel instance
125306c3fb27SDimitry Andric // alignment padding
125406c3fb27SDimitry Andric // dynamic lds variables
125506c3fb27SDimitry Andric //}
125606c3fb27SDimitry Andric
125706c3fb27SDimitry Andric const bool AllocateModuleScopeStruct =
125806c3fb27SDimitry Andric MaybeModuleScopeStruct &&
125906c3fb27SDimitry Andric KernelsThatAllocateModuleLDS.contains(&Func);
126006c3fb27SDimitry Andric
126106c3fb27SDimitry Andric auto Replacement = KernelToReplacement.find(&Func);
126206c3fb27SDimitry Andric const bool AllocateKernelScopeStruct =
126306c3fb27SDimitry Andric Replacement != KernelToReplacement.end();
126406c3fb27SDimitry Andric
126506c3fb27SDimitry Andric const bool AllocateDynamicVariable =
126606c3fb27SDimitry Andric KernelToCreatedDynamicLDS.contains(&Func);
126706c3fb27SDimitry Andric
126806c3fb27SDimitry Andric uint32_t Offset = 0;
126906c3fb27SDimitry Andric
127006c3fb27SDimitry Andric if (AllocateModuleScopeStruct) {
127106c3fb27SDimitry Andric // Allocated at zero, recorded once on construction, not once per
127206c3fb27SDimitry Andric // kernel
127306c3fb27SDimitry Andric Offset += DL.getTypeAllocSize(MaybeModuleScopeStruct->getValueType());
127406c3fb27SDimitry Andric }
127506c3fb27SDimitry Andric
127606c3fb27SDimitry Andric if (AllocateKernelScopeStruct) {
127706c3fb27SDimitry Andric GlobalVariable *KernelStruct = Replacement->second.SGV;
127806c3fb27SDimitry Andric Offset = alignTo(Offset, AMDGPU::getAlign(DL, KernelStruct));
127906c3fb27SDimitry Andric recordLDSAbsoluteAddress(&M, KernelStruct, Offset);
128006c3fb27SDimitry Andric Offset += DL.getTypeAllocSize(KernelStruct->getValueType());
128106c3fb27SDimitry Andric }
128206c3fb27SDimitry Andric
128306c3fb27SDimitry Andric // If there is dynamic allocation, the alignment needed is included in
128406c3fb27SDimitry Andric // the static frame size. There may be no reference to the dynamic
128506c3fb27SDimitry Andric // variable in the kernel itself, so without including it here, that
128606c3fb27SDimitry Andric // alignment padding could be missed.
128706c3fb27SDimitry Andric if (AllocateDynamicVariable) {
128806c3fb27SDimitry Andric GlobalVariable *DynamicVariable = KernelToCreatedDynamicLDS[&Func];
128906c3fb27SDimitry Andric Offset = alignTo(Offset, AMDGPU::getAlign(DL, DynamicVariable));
129006c3fb27SDimitry Andric recordLDSAbsoluteAddress(&M, DynamicVariable, Offset);
129106c3fb27SDimitry Andric }
129206c3fb27SDimitry Andric
129306c3fb27SDimitry Andric if (Offset != 0) {
12945f757f3fSDimitry Andric (void)TM; // TODO: Account for target maximum LDS
129506c3fb27SDimitry Andric std::string Buffer;
129606c3fb27SDimitry Andric raw_string_ostream SS{Buffer};
129706c3fb27SDimitry Andric SS << format("%u", Offset);
129806c3fb27SDimitry Andric
129906c3fb27SDimitry Andric // Instead of explictly marking kernels that access dynamic variables
130006c3fb27SDimitry Andric // using special case metadata, annotate with min-lds == max-lds, i.e.
130106c3fb27SDimitry Andric // that there is no more space available for allocating more static
130206c3fb27SDimitry Andric // LDS variables. That is the right condition to prevent allocating
130306c3fb27SDimitry Andric // more variables which would collide with the addresses assigned to
130406c3fb27SDimitry Andric // dynamic variables.
130506c3fb27SDimitry Andric if (AllocateDynamicVariable)
130606c3fb27SDimitry Andric SS << format(",%u", Offset);
130706c3fb27SDimitry Andric
130806c3fb27SDimitry Andric Func.addFnAttr("amdgpu-lds-size", Buffer);
130906c3fb27SDimitry Andric }
131006c3fb27SDimitry Andric }
131106c3fb27SDimitry Andric }
131206c3fb27SDimitry Andric
1313bdd1243dSDimitry Andric for (auto &GV : make_early_inc_range(M.globals()))
1314bdd1243dSDimitry Andric if (AMDGPU::isLDSVariableToLower(GV)) {
1315bdd1243dSDimitry Andric // probably want to remove from used lists
1316bdd1243dSDimitry Andric GV.removeDeadConstantUsers();
1317bdd1243dSDimitry Andric if (GV.use_empty())
1318bdd1243dSDimitry Andric GV.eraseFromParent();
1319fe6060f1SDimitry Andric }
1320fe6060f1SDimitry Andric
1321fe6060f1SDimitry Andric return Changed;
1322fe6060f1SDimitry Andric }
1323fe6060f1SDimitry Andric
1324fe6060f1SDimitry Andric private:
1325fe6060f1SDimitry Andric // Increase the alignment of LDS globals if necessary to maximise the chance
1326fe6060f1SDimitry Andric // that we can use aligned LDS instructions to access them.
superAlignLDSGlobals(Module & M)13270eae32dcSDimitry Andric static bool superAlignLDSGlobals(Module &M) {
13280eae32dcSDimitry Andric const DataLayout &DL = M.getDataLayout();
13290eae32dcSDimitry Andric bool Changed = false;
13300eae32dcSDimitry Andric if (!SuperAlignLDSGlobals) {
13310eae32dcSDimitry Andric return Changed;
13320eae32dcSDimitry Andric }
13330eae32dcSDimitry Andric
13340eae32dcSDimitry Andric for (auto &GV : M.globals()) {
13350eae32dcSDimitry Andric if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {
13360eae32dcSDimitry Andric // Only changing alignment of LDS variables
13370eae32dcSDimitry Andric continue;
13380eae32dcSDimitry Andric }
13390eae32dcSDimitry Andric if (!GV.hasInitializer()) {
13400eae32dcSDimitry Andric // cuda/hip extern __shared__ variable, leave alignment alone
13410eae32dcSDimitry Andric continue;
13420eae32dcSDimitry Andric }
13430eae32dcSDimitry Andric
13440eae32dcSDimitry Andric Align Alignment = AMDGPU::getAlign(DL, &GV);
13450eae32dcSDimitry Andric TypeSize GVSize = DL.getTypeAllocSize(GV.getValueType());
1346fe6060f1SDimitry Andric
1347fe6060f1SDimitry Andric if (GVSize > 8) {
1348fe6060f1SDimitry Andric // We might want to use a b96 or b128 load/store
1349fe6060f1SDimitry Andric Alignment = std::max(Alignment, Align(16));
1350fe6060f1SDimitry Andric } else if (GVSize > 4) {
1351fe6060f1SDimitry Andric // We might want to use a b64 load/store
1352fe6060f1SDimitry Andric Alignment = std::max(Alignment, Align(8));
1353fe6060f1SDimitry Andric } else if (GVSize > 2) {
1354fe6060f1SDimitry Andric // We might want to use a b32 load/store
1355fe6060f1SDimitry Andric Alignment = std::max(Alignment, Align(4));
1356fe6060f1SDimitry Andric } else if (GVSize > 1) {
1357fe6060f1SDimitry Andric // We might want to use a b16 load/store
1358fe6060f1SDimitry Andric Alignment = std::max(Alignment, Align(2));
1359fe6060f1SDimitry Andric }
1360fe6060f1SDimitry Andric
13610eae32dcSDimitry Andric if (Alignment != AMDGPU::getAlign(DL, &GV)) {
13620eae32dcSDimitry Andric Changed = true;
13630eae32dcSDimitry Andric GV.setAlignment(Alignment);
1364fe6060f1SDimitry Andric }
1365fe6060f1SDimitry Andric }
13660eae32dcSDimitry Andric return Changed;
13670eae32dcSDimitry Andric }
13680eae32dcSDimitry Andric
createLDSVariableReplacement(Module & M,std::string VarName,DenseSet<GlobalVariable * > const & LDSVarsToTransform)1369bdd1243dSDimitry Andric static LDSVariableReplacement createLDSVariableReplacement(
1370972a253aSDimitry Andric Module &M, std::string VarName,
1371bdd1243dSDimitry Andric DenseSet<GlobalVariable *> const &LDSVarsToTransform) {
1372972a253aSDimitry Andric // Create a struct instance containing LDSVarsToTransform and map from those
1373972a253aSDimitry Andric // variables to ConstantExprGEP
1374972a253aSDimitry Andric // Variables may be introduced to meet alignment requirements. No aliasing
1375972a253aSDimitry Andric // metadata is useful for these as they have no uses. Erased before return.
1376972a253aSDimitry Andric
13770eae32dcSDimitry Andric LLVMContext &Ctx = M.getContext();
13780eae32dcSDimitry Andric const DataLayout &DL = M.getDataLayout();
1379972a253aSDimitry Andric assert(!LDSVarsToTransform.empty());
1380fe6060f1SDimitry Andric
1381fe6060f1SDimitry Andric SmallVector<OptimizedStructLayoutField, 8> LayoutFields;
1382fcaf7f86SDimitry Andric LayoutFields.reserve(LDSVarsToTransform.size());
1383bdd1243dSDimitry Andric {
1384bdd1243dSDimitry Andric // The order of fields in this struct depends on the order of
1385bdd1243dSDimitry Andric // varables in the argument which varies when changing how they
1386bdd1243dSDimitry Andric // are identified, leading to spurious test breakage.
138706c3fb27SDimitry Andric auto Sorted = sortByName(std::vector<GlobalVariable *>(
138806c3fb27SDimitry Andric LDSVarsToTransform.begin(), LDSVarsToTransform.end()));
138906c3fb27SDimitry Andric
1390bdd1243dSDimitry Andric for (GlobalVariable *GV : Sorted) {
1391bdd1243dSDimitry Andric OptimizedStructLayoutField F(GV,
1392bdd1243dSDimitry Andric DL.getTypeAllocSize(GV->getValueType()),
1393fe6060f1SDimitry Andric AMDGPU::getAlign(DL, GV));
1394fe6060f1SDimitry Andric LayoutFields.emplace_back(F);
1395fe6060f1SDimitry Andric }
1396bdd1243dSDimitry Andric }
1397fe6060f1SDimitry Andric
1398fe6060f1SDimitry Andric performOptimizedStructLayout(LayoutFields);
1399fe6060f1SDimitry Andric
1400fe6060f1SDimitry Andric std::vector<GlobalVariable *> LocalVars;
1401972a253aSDimitry Andric BitVector IsPaddingField;
1402fcaf7f86SDimitry Andric LocalVars.reserve(LDSVarsToTransform.size()); // will be at least this large
1403972a253aSDimitry Andric IsPaddingField.reserve(LDSVarsToTransform.size());
1404fe6060f1SDimitry Andric {
1405fe6060f1SDimitry Andric uint64_t CurrentOffset = 0;
1406fe6060f1SDimitry Andric for (size_t I = 0; I < LayoutFields.size(); I++) {
1407fe6060f1SDimitry Andric GlobalVariable *FGV = static_cast<GlobalVariable *>(
1408fe6060f1SDimitry Andric const_cast<void *>(LayoutFields[I].Id));
1409fe6060f1SDimitry Andric Align DataAlign = LayoutFields[I].Alignment;
1410fe6060f1SDimitry Andric
1411fe6060f1SDimitry Andric uint64_t DataAlignV = DataAlign.value();
1412fe6060f1SDimitry Andric if (uint64_t Rem = CurrentOffset % DataAlignV) {
1413fe6060f1SDimitry Andric uint64_t Padding = DataAlignV - Rem;
1414fe6060f1SDimitry Andric
1415fe6060f1SDimitry Andric // Append an array of padding bytes to meet alignment requested
1416fe6060f1SDimitry Andric // Note (o + (a - (o % a)) ) % a == 0
1417fe6060f1SDimitry Andric // (offset + Padding ) % align == 0
1418fe6060f1SDimitry Andric
1419fe6060f1SDimitry Andric Type *ATy = ArrayType::get(Type::getInt8Ty(Ctx), Padding);
1420fe6060f1SDimitry Andric LocalVars.push_back(new GlobalVariable(
14215f757f3fSDimitry Andric M, ATy, false, GlobalValue::InternalLinkage,
14225f757f3fSDimitry Andric PoisonValue::get(ATy), "", nullptr, GlobalValue::NotThreadLocal,
14235f757f3fSDimitry Andric AMDGPUAS::LOCAL_ADDRESS, false));
1424972a253aSDimitry Andric IsPaddingField.push_back(true);
1425fe6060f1SDimitry Andric CurrentOffset += Padding;
1426fe6060f1SDimitry Andric }
1427fe6060f1SDimitry Andric
1428fe6060f1SDimitry Andric LocalVars.push_back(FGV);
1429972a253aSDimitry Andric IsPaddingField.push_back(false);
1430fe6060f1SDimitry Andric CurrentOffset += LayoutFields[I].Size;
1431fe6060f1SDimitry Andric }
1432fe6060f1SDimitry Andric }
1433fe6060f1SDimitry Andric
1434fe6060f1SDimitry Andric std::vector<Type *> LocalVarTypes;
1435fe6060f1SDimitry Andric LocalVarTypes.reserve(LocalVars.size());
1436fe6060f1SDimitry Andric std::transform(
1437fe6060f1SDimitry Andric LocalVars.cbegin(), LocalVars.cend(), std::back_inserter(LocalVarTypes),
1438fe6060f1SDimitry Andric [](const GlobalVariable *V) -> Type * { return V->getValueType(); });
1439fe6060f1SDimitry Andric
1440fe6060f1SDimitry Andric StructType *LDSTy = StructType::create(Ctx, LocalVarTypes, VarName + ".t");
1441fe6060f1SDimitry Andric
1442bdd1243dSDimitry Andric Align StructAlign = AMDGPU::getAlign(DL, LocalVars[0]);
1443fe6060f1SDimitry Andric
1444fe6060f1SDimitry Andric GlobalVariable *SGV = new GlobalVariable(
14455f757f3fSDimitry Andric M, LDSTy, false, GlobalValue::InternalLinkage, PoisonValue::get(LDSTy),
1446fe6060f1SDimitry Andric VarName, nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS,
1447fe6060f1SDimitry Andric false);
1448fe6060f1SDimitry Andric SGV->setAlignment(StructAlign);
1449972a253aSDimitry Andric
1450972a253aSDimitry Andric DenseMap<GlobalVariable *, Constant *> Map;
1451972a253aSDimitry Andric Type *I32 = Type::getInt32Ty(Ctx);
1452972a253aSDimitry Andric for (size_t I = 0; I < LocalVars.size(); I++) {
1453972a253aSDimitry Andric GlobalVariable *GV = LocalVars[I];
1454972a253aSDimitry Andric Constant *GEPIdx[] = {ConstantInt::get(I32, 0), ConstantInt::get(I32, I)};
1455972a253aSDimitry Andric Constant *GEP = ConstantExpr::getGetElementPtr(LDSTy, SGV, GEPIdx, true);
1456972a253aSDimitry Andric if (IsPaddingField[I]) {
1457972a253aSDimitry Andric assert(GV->use_empty());
1458972a253aSDimitry Andric GV->eraseFromParent();
1459972a253aSDimitry Andric } else {
1460972a253aSDimitry Andric Map[GV] = GEP;
1461972a253aSDimitry Andric }
1462972a253aSDimitry Andric }
1463972a253aSDimitry Andric assert(Map.size() == LDSVarsToTransform.size());
1464972a253aSDimitry Andric return {SGV, std::move(Map)};
1465fe6060f1SDimitry Andric }
1466fe6060f1SDimitry Andric
1467972a253aSDimitry Andric template <typename PredicateTy>
replaceLDSVariablesWithStruct(Module & M,DenseSet<GlobalVariable * > const & LDSVarsToTransformArg,const LDSVariableReplacement & Replacement,PredicateTy Predicate)146806c3fb27SDimitry Andric static void replaceLDSVariablesWithStruct(
1469bdd1243dSDimitry Andric Module &M, DenseSet<GlobalVariable *> const &LDSVarsToTransformArg,
147006c3fb27SDimitry Andric const LDSVariableReplacement &Replacement, PredicateTy Predicate) {
1471972a253aSDimitry Andric LLVMContext &Ctx = M.getContext();
1472972a253aSDimitry Andric const DataLayout &DL = M.getDataLayout();
1473fe6060f1SDimitry Andric
1474bdd1243dSDimitry Andric // A hack... we need to insert the aliasing info in a predictable order for
1475bdd1243dSDimitry Andric // lit tests. Would like to have them in a stable order already, ideally the
1476bdd1243dSDimitry Andric // same order they get allocated, which might mean an ordered set container
147706c3fb27SDimitry Andric auto LDSVarsToTransform = sortByName(std::vector<GlobalVariable *>(
147806c3fb27SDimitry Andric LDSVarsToTransformArg.begin(), LDSVarsToTransformArg.end()));
1479bdd1243dSDimitry Andric
1480349cc55cSDimitry Andric // Create alias.scope and their lists. Each field in the new structure
1481349cc55cSDimitry Andric // does not alias with all other fields.
1482349cc55cSDimitry Andric SmallVector<MDNode *> AliasScopes;
1483349cc55cSDimitry Andric SmallVector<Metadata *> NoAliasList;
1484972a253aSDimitry Andric const size_t NumberVars = LDSVarsToTransform.size();
1485972a253aSDimitry Andric if (NumberVars > 1) {
1486349cc55cSDimitry Andric MDBuilder MDB(Ctx);
1487972a253aSDimitry Andric AliasScopes.reserve(NumberVars);
1488349cc55cSDimitry Andric MDNode *Domain = MDB.createAnonymousAliasScopeDomain();
1489972a253aSDimitry Andric for (size_t I = 0; I < NumberVars; I++) {
1490349cc55cSDimitry Andric MDNode *Scope = MDB.createAnonymousAliasScope(Domain);
1491349cc55cSDimitry Andric AliasScopes.push_back(Scope);
1492349cc55cSDimitry Andric }
1493349cc55cSDimitry Andric NoAliasList.append(&AliasScopes[1], AliasScopes.end());
1494349cc55cSDimitry Andric }
1495349cc55cSDimitry Andric
1496972a253aSDimitry Andric // Replace uses of ith variable with a constantexpr to the corresponding
1497972a253aSDimitry Andric // field of the instance that will be allocated by AMDGPUMachineFunction
1498972a253aSDimitry Andric for (size_t I = 0; I < NumberVars; I++) {
1499972a253aSDimitry Andric GlobalVariable *GV = LDSVarsToTransform[I];
150006c3fb27SDimitry Andric Constant *GEP = Replacement.LDSVarsToConstantGEP.at(GV);
1501fe6060f1SDimitry Andric
1502972a253aSDimitry Andric GV->replaceUsesWithIf(GEP, Predicate);
1503fe6060f1SDimitry Andric
1504972a253aSDimitry Andric APInt APOff(DL.getIndexTypeSizeInBits(GEP->getType()), 0);
1505972a253aSDimitry Andric GEP->stripAndAccumulateInBoundsConstantOffsets(DL, APOff);
1506972a253aSDimitry Andric uint64_t Offset = APOff.getZExtValue();
1507972a253aSDimitry Andric
1508bdd1243dSDimitry Andric Align A =
1509bdd1243dSDimitry Andric commonAlignment(Replacement.SGV->getAlign().valueOrOne(), Offset);
1510349cc55cSDimitry Andric
1511349cc55cSDimitry Andric if (I)
1512349cc55cSDimitry Andric NoAliasList[I - 1] = AliasScopes[I - 1];
1513349cc55cSDimitry Andric MDNode *NoAlias =
1514349cc55cSDimitry Andric NoAliasList.empty() ? nullptr : MDNode::get(Ctx, NoAliasList);
1515349cc55cSDimitry Andric MDNode *AliasScope =
1516349cc55cSDimitry Andric AliasScopes.empty() ? nullptr : MDNode::get(Ctx, {AliasScopes[I]});
1517349cc55cSDimitry Andric
1518349cc55cSDimitry Andric refineUsesAlignmentAndAA(GEP, A, DL, AliasScope, NoAlias);
1519fe6060f1SDimitry Andric }
1520fe6060f1SDimitry Andric }
1521fe6060f1SDimitry Andric
refineUsesAlignmentAndAA(Value * Ptr,Align A,const DataLayout & DL,MDNode * AliasScope,MDNode * NoAlias,unsigned MaxDepth=5)152206c3fb27SDimitry Andric static void refineUsesAlignmentAndAA(Value *Ptr, Align A,
152306c3fb27SDimitry Andric const DataLayout &DL, MDNode *AliasScope,
152406c3fb27SDimitry Andric MDNode *NoAlias, unsigned MaxDepth = 5) {
1525349cc55cSDimitry Andric if (!MaxDepth || (A == 1 && !AliasScope))
1526fe6060f1SDimitry Andric return;
1527fe6060f1SDimitry Andric
1528fe6060f1SDimitry Andric for (User *U : Ptr->users()) {
1529349cc55cSDimitry Andric if (auto *I = dyn_cast<Instruction>(U)) {
1530349cc55cSDimitry Andric if (AliasScope && I->mayReadOrWriteMemory()) {
1531349cc55cSDimitry Andric MDNode *AS = I->getMetadata(LLVMContext::MD_alias_scope);
1532349cc55cSDimitry Andric AS = (AS ? MDNode::getMostGenericAliasScope(AS, AliasScope)
1533349cc55cSDimitry Andric : AliasScope);
1534349cc55cSDimitry Andric I->setMetadata(LLVMContext::MD_alias_scope, AS);
1535349cc55cSDimitry Andric
1536349cc55cSDimitry Andric MDNode *NA = I->getMetadata(LLVMContext::MD_noalias);
1537349cc55cSDimitry Andric NA = (NA ? MDNode::intersect(NA, NoAlias) : NoAlias);
1538349cc55cSDimitry Andric I->setMetadata(LLVMContext::MD_noalias, NA);
1539349cc55cSDimitry Andric }
1540349cc55cSDimitry Andric }
1541349cc55cSDimitry Andric
1542fe6060f1SDimitry Andric if (auto *LI = dyn_cast<LoadInst>(U)) {
1543fe6060f1SDimitry Andric LI->setAlignment(std::max(A, LI->getAlign()));
1544fe6060f1SDimitry Andric continue;
1545fe6060f1SDimitry Andric }
1546fe6060f1SDimitry Andric if (auto *SI = dyn_cast<StoreInst>(U)) {
1547fe6060f1SDimitry Andric if (SI->getPointerOperand() == Ptr)
1548fe6060f1SDimitry Andric SI->setAlignment(std::max(A, SI->getAlign()));
1549fe6060f1SDimitry Andric continue;
1550fe6060f1SDimitry Andric }
1551fe6060f1SDimitry Andric if (auto *AI = dyn_cast<AtomicRMWInst>(U)) {
1552fe6060f1SDimitry Andric // None of atomicrmw operations can work on pointers, but let's
1553fe6060f1SDimitry Andric // check it anyway in case it will or we will process ConstantExpr.
1554fe6060f1SDimitry Andric if (AI->getPointerOperand() == Ptr)
1555fe6060f1SDimitry Andric AI->setAlignment(std::max(A, AI->getAlign()));
1556fe6060f1SDimitry Andric continue;
1557fe6060f1SDimitry Andric }
1558fe6060f1SDimitry Andric if (auto *AI = dyn_cast<AtomicCmpXchgInst>(U)) {
1559fe6060f1SDimitry Andric if (AI->getPointerOperand() == Ptr)
1560fe6060f1SDimitry Andric AI->setAlignment(std::max(A, AI->getAlign()));
1561fe6060f1SDimitry Andric continue;
1562fe6060f1SDimitry Andric }
1563fe6060f1SDimitry Andric if (auto *GEP = dyn_cast<GetElementPtrInst>(U)) {
1564fe6060f1SDimitry Andric unsigned BitWidth = DL.getIndexTypeSizeInBits(GEP->getType());
1565fe6060f1SDimitry Andric APInt Off(BitWidth, 0);
1566349cc55cSDimitry Andric if (GEP->getPointerOperand() == Ptr) {
1567349cc55cSDimitry Andric Align GA;
1568349cc55cSDimitry Andric if (GEP->accumulateConstantOffset(DL, Off))
1569349cc55cSDimitry Andric GA = commonAlignment(A, Off.getLimitedValue());
1570349cc55cSDimitry Andric refineUsesAlignmentAndAA(GEP, GA, DL, AliasScope, NoAlias,
1571349cc55cSDimitry Andric MaxDepth - 1);
1572fe6060f1SDimitry Andric }
1573fe6060f1SDimitry Andric continue;
1574fe6060f1SDimitry Andric }
1575fe6060f1SDimitry Andric if (auto *I = dyn_cast<Instruction>(U)) {
1576fe6060f1SDimitry Andric if (I->getOpcode() == Instruction::BitCast ||
1577fe6060f1SDimitry Andric I->getOpcode() == Instruction::AddrSpaceCast)
1578349cc55cSDimitry Andric refineUsesAlignmentAndAA(I, A, DL, AliasScope, NoAlias, MaxDepth - 1);
1579fe6060f1SDimitry Andric }
1580fe6060f1SDimitry Andric }
1581fe6060f1SDimitry Andric }
1582fe6060f1SDimitry Andric };
1583fe6060f1SDimitry Andric
15845f757f3fSDimitry Andric class AMDGPULowerModuleLDSLegacy : public ModulePass {
15855f757f3fSDimitry Andric public:
15865f757f3fSDimitry Andric const AMDGPUTargetMachine *TM;
15875f757f3fSDimitry Andric static char ID;
15885f757f3fSDimitry Andric
AMDGPULowerModuleLDSLegacy(const AMDGPUTargetMachine * TM_=nullptr)15895f757f3fSDimitry Andric AMDGPULowerModuleLDSLegacy(const AMDGPUTargetMachine *TM_ = nullptr)
15905f757f3fSDimitry Andric : ModulePass(ID), TM(TM_) {
15915f757f3fSDimitry Andric initializeAMDGPULowerModuleLDSLegacyPass(*PassRegistry::getPassRegistry());
15925f757f3fSDimitry Andric }
15935f757f3fSDimitry Andric
getAnalysisUsage(AnalysisUsage & AU) const15945f757f3fSDimitry Andric void getAnalysisUsage(AnalysisUsage &AU) const override {
15955f757f3fSDimitry Andric if (!TM)
15965f757f3fSDimitry Andric AU.addRequired<TargetPassConfig>();
15975f757f3fSDimitry Andric }
15985f757f3fSDimitry Andric
runOnModule(Module & M)15995f757f3fSDimitry Andric bool runOnModule(Module &M) override {
16005f757f3fSDimitry Andric if (!TM) {
16015f757f3fSDimitry Andric auto &TPC = getAnalysis<TargetPassConfig>();
16025f757f3fSDimitry Andric TM = &TPC.getTM<AMDGPUTargetMachine>();
16035f757f3fSDimitry Andric }
16045f757f3fSDimitry Andric
16055f757f3fSDimitry Andric return AMDGPULowerModuleLDS(*TM).runOnModule(M);
16065f757f3fSDimitry Andric }
16075f757f3fSDimitry Andric };
16085f757f3fSDimitry Andric
1609fe6060f1SDimitry Andric } // namespace
16105f757f3fSDimitry Andric char AMDGPULowerModuleLDSLegacy::ID = 0;
1611fe6060f1SDimitry Andric
16125f757f3fSDimitry Andric char &llvm::AMDGPULowerModuleLDSLegacyPassID = AMDGPULowerModuleLDSLegacy::ID;
1613fe6060f1SDimitry Andric
16145f757f3fSDimitry Andric INITIALIZE_PASS_BEGIN(AMDGPULowerModuleLDSLegacy, DEBUG_TYPE,
16155f757f3fSDimitry Andric "Lower uses of LDS variables from non-kernel functions",
16165f757f3fSDimitry Andric false, false)
INITIALIZE_PASS_DEPENDENCY(TargetPassConfig)16175f757f3fSDimitry Andric INITIALIZE_PASS_DEPENDENCY(TargetPassConfig)
16185f757f3fSDimitry Andric INITIALIZE_PASS_END(AMDGPULowerModuleLDSLegacy, DEBUG_TYPE,
16195f757f3fSDimitry Andric "Lower uses of LDS variables from non-kernel functions",
16205f757f3fSDimitry Andric false, false)
1621fe6060f1SDimitry Andric
16225f757f3fSDimitry Andric ModulePass *
16235f757f3fSDimitry Andric llvm::createAMDGPULowerModuleLDSLegacyPass(const AMDGPUTargetMachine *TM) {
16245f757f3fSDimitry Andric return new AMDGPULowerModuleLDSLegacy(TM);
1625fe6060f1SDimitry Andric }
1626fe6060f1SDimitry Andric
run(Module & M,ModuleAnalysisManager &)1627fe6060f1SDimitry Andric PreservedAnalyses AMDGPULowerModuleLDSPass::run(Module &M,
1628fe6060f1SDimitry Andric ModuleAnalysisManager &) {
16295f757f3fSDimitry Andric return AMDGPULowerModuleLDS(TM).runOnModule(M) ? PreservedAnalyses::none()
1630fe6060f1SDimitry Andric : PreservedAnalyses::all();
1631fe6060f1SDimitry Andric }
1632