1201342Snyan//===-- AMDGPULowerModuleLDSPass.cpp ------------------------------*- C++ -*-=// 2201342Snyan// 3201342Snyan// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4201342Snyan// See https://llvm.org/LICENSE.txt for license information. 5201342Snyan// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6201342Snyan// 7201342Snyan//===----------------------------------------------------------------------===// 8201342Snyan// 9201342Snyan// This pass eliminates local data store, LDS, uses from non-kernel functions. 10201342Snyan// LDS is contiguous memory allocated per kernel execution. 11201342Snyan// 12201342Snyan// Background. 13201342Snyan// 14201342Snyan// The programming model is global variables, or equivalently function local 15201342Snyan// static variables, accessible from kernels or other functions. For uses from 16201342Snyan// kernels this is straightforward - assign an integer to the kernel for the 17201342Snyan// memory required by all the variables combined, allocate them within that. 18201342Snyan// For uses from functions there are performance tradeoffs to choose between. 19201342Snyan// 20201342Snyan// This model means the GPU runtime can specify the amount of memory allocated. 21201342Snyan// If this is more than the kernel assumed, the excess can be made available 22201342Snyan// using a language specific feature, which IR represents as a variable with 23201342Snyan// no initializer. This feature is referred to here as "Dynamic LDS" and is 24201342Snyan// lowered slightly differently to the normal case. 25201342Snyan// 26201342Snyan// Consequences of this GPU feature: 27201342Snyan// - memory is limited and exceeding it halts compilation 28201342Snyan// - a global accessed by one kernel exists independent of other kernels 29201342Snyan// - a global exists independent of simultaneous execution of the same kernel 30201342Snyan// - the address of the global may be different from different kernels as they 31201342Snyan// do not alias, which permits only allocating variables they use 32201342Snyan// - if the address is allowed to differ, functions need help to find it 33201342Snyan// 34201342Snyan// Uses from kernels are implemented here by grouping them in a per-kernel 35201342Snyan// struct instance. This duplicates the variables, accurately modelling their 36201342Snyan// aliasing properties relative to a single global representation. It also 37201342Snyan// permits control over alignment via padding. 38201342Snyan// 39201342Snyan// Uses from functions are more complicated and the primary purpose of this 40201342Snyan// IR pass. Several different lowering are chosen between to meet requirements 41201342Snyan// to avoid allocating any LDS where it is not necessary, as that impacts 42201342Snyan// occupancy and may fail the compilation, while not imposing overhead on a 43201342Snyan// feature whose primary advantage over global memory is performance. The basic 44201342Snyan// design goal is to avoid one kernel imposing overhead on another. 45201342Snyan// 46201342Snyan// Implementation. 47201342Snyan// 48201342Snyan// LDS variables with constant annotation or non-undef initializer are passed 49201342Snyan// through unchanged for simplification or error diagnostics in later passes. 50201342Snyan// Non-undef initializers are not yet implemented for LDS. 51201342Snyan// 52201342Snyan// LDS variables that are always allocated at the same address can be found 53201342Snyan// by lookup at that address. Otherwise runtime information/cost is required. 54201342Snyan// 55201342Snyan// The simplest strategy possible is to group all LDS variables in a single 56201342Snyan// struct and allocate that struct in every kernel such that the original 57201342Snyan// variables are always at the same address. LDS is however a limited resource 58201342Snyan// so this strategy is unusable in practice. It is not implemented here. 59201342Snyan// 60201342Snyan// Strategy | Precise allocation | Zero runtime cost | General purpose | 61201342Snyan// --------+--------------------+-------------------+-----------------+ 62201342Snyan// Module | No | Yes | Yes | 63201342Snyan// Table | Yes | No | Yes | 64201342Snyan// Kernel | Yes | Yes | No | 65201342Snyan// Hybrid | Yes | Partial | Yes | 66201342Snyan// 67201342Snyan// "Module" spends LDS memory to save cycles. "Table" spends cycles and global 68201342Snyan// memory to save LDS. "Kernel" is as fast as kernel allocation but only works 69201342Snyan// for variables that are known reachable from a single kernel. "Hybrid" picks 70201342Snyan// between all three. When forced to choose between LDS and cycles we minimise 71201342Snyan// LDS use. 72201342Snyan 73201342Snyan// The "module" lowering implemented here finds LDS variables which are used by 74201342Snyan// non-kernel functions and creates a new struct with a field for each of those 75201342Snyan// LDS variables. Variables that are only used from kernels are excluded. 76201342Snyan// 77201342Snyan// The "table" lowering implemented here has three components. 78201342Snyan// First kernels are assigned a unique integer identifier which is available in 79226506Sdes// functions it calls through the intrinsic amdgcn_lds_kernel_id. The integer 80226506Sdes// is passed through a specific SGPR, thus works with indirect calls. 81201342Snyan// Second, each kernel allocates LDS variables independent of other kernels and 82201342Snyan// writes the addresses it chose for each variable into an array in consistent 83201342Snyan// order. If the kernel does not allocate a given variable, it writes undef to 84201342Snyan// the corresponding array location. These arrays are written to a constant 85201342Snyan// table in the order matching the kernel unique integer identifier. 86201342Snyan// Third, uses from non-kernel functions are replaced with a table lookup using 87201342Snyan// the intrinsic function to find the address of the variable. 88201342Snyan// 89201342Snyan// "Kernel" lowering is only applicable for variables that are unambiguously 90201342Snyan// reachable from exactly one kernel. For those cases, accesses to the variable 91201342Snyan// can be lowered to ConstantExpr address of a struct instance specific to that 92201342Snyan// one kernel. This is zero cost in space and in compute. It will raise a fatal 93201342Snyan// error on any variable that might be reachable from multiple kernels and is 94201342Snyan// thus most easily used as part of the hybrid lowering strategy. 95201342Snyan// 96201342Snyan// Hybrid lowering is a mixture of the above. It uses the zero cost kernel 97201342Snyan// lowering where it can. It lowers the variable accessed by the greatest 98201342Snyan// number of kernels using the module strategy as that is free for the first 99201342Snyan// variable. Any futher variables that can be lowered with the module strategy 100201342Snyan// without incurring LDS memory overhead are. The remaining ones are lowered 101201342Snyan// via table. 102201342Snyan// 103201342Snyan// Consequences 104201342Snyan// - No heuristics or user controlled magic numbers, hybrid is the right choice 105201342Snyan// - Kernels that don't use functions (or have had them all inlined) are not 106201342Snyan// affected by any lowering for kernels that do. 107201342Snyan// - Kernels that don't make indirect function calls are not affected by those 108201342Snyan// that do. 109201342Snyan// - Variables which are used by lots of kernels, e.g. those injected by a 110201342Snyan// language runtime in most kernels, are expected to have no overhead 111201342Snyan// - Implementations that instantiate templates per-kernel where those templates 112201342Snyan// use LDS are expected to hit the "Kernel" lowering strategy 113201342Snyan// - The runtime properties impose a cost in compiler implementation complexity 114201342Snyan// 115201342Snyan// Dynamic LDS implementation 116201342Snyan// Dynamic LDS is lowered similarly to the "table" strategy above and uses the 117201342Snyan// same intrinsic to identify which kernel is at the root of the dynamic call 118201342Snyan// graph. This relies on the specified behaviour that all dynamic LDS variables 119201342Snyan// alias one another, i.e. are at the same address, with respect to a given 120201342Snyan// kernel. Therefore this pass creates new dynamic LDS variables for each kernel 121201342Snyan// that allocates any dynamic LDS and builds a table of addresses out of those. 122201342Snyan// The AMDGPUPromoteAlloca pass skips kernels that use dynamic LDS. 123201342Snyan// The corresponding optimisation for "kernel" lowering where the table lookup 124201342Snyan// is elided is not implemented. 125201342Snyan// 126201342Snyan// 127201342Snyan// Implementation notes / limitations 128201342Snyan// A single LDS global variable represents an instance per kernel that can reach 129239063Snyan// said variables. This pass essentially specialises said variables per kernel. 130239063Snyan// Handling ConstantExpr during the pass complicated this significantly so now 131201342Snyan// all ConstantExpr uses of LDS variables are expanded to instructions. This 132201342Snyan// may need amending when implementing non-undef initialisers. 133232784Snyan// 134239063Snyan// Lowering is split between this IR pass and the back end. This pass chooses 135239063Snyan// where given variables should be allocated and marks them with metadata, 136201342Snyan// MD_absolute_symbol. The backend places the variables in coincidentally the 137201342Snyan// same location and raises a fatal error if something has gone awry. This works 138219960Snyan// in practice because the only pass between this one and the backend that 139201342Snyan// changes LDS is PromoteAlloca and the changes it makes do not conflict. 140201342Snyan// 141201342Snyan// Addresses are written to constant global arrays based on the same metadata. 142201342Snyan// 143201342Snyan// The backend lowers LDS variables in the order of traversal of the function. 144201342Snyan// This is at odds with the deterministic layout required. The workaround is to 145201342Snyan// allocate the fixed-address variables immediately upon starting the function 146201342Snyan// where they can be placed as intended. This requires a means of mapping from 147201342Snyan// the function to the variables that it allocates. For the module scope lds, 148201342Snyan// this is via metadata indicating whether the variable is not required. If a 149201342Snyan// pass deletes that metadata, a fatal error on disagreement with the absolute 150220685Snyan// symbol metadata will occur. For kernel scope and dynamic, this is by _name_ 151201342Snyan// correspondence between the function and the variable. It requires the 152201342Snyan// kernel to have a name (which is only a limitation for tests in practice) and 153201342Snyan// for nothing to rename the corresponding symbols. This is a hazard if the pass 154201342Snyan// is run multiple times during debugging. Alternative schemes considered all 155201342Snyan// involve bespoke metadata. 156201342Snyan// 157201342Snyan// If the name correspondence can be replaced, multiple distinct kernels that 158201342Snyan// have the same memory layout can map to the same kernel id (as the address 159201342Snyan// itself is handled by the absolute symbol metadata) and that will allow more 160201342Snyan// uses of the "kernel" style faster lowering and reduce the size of the lookup 161201342Snyan// tables. 162201342Snyan// 163201342Snyan// There is a test that checks this does not fire for a graphics shader. This 164201342Snyan// lowering is expected to work for graphics if the isKernel test is changed. 165201342Snyan// 166201342Snyan// The current markUsedByKernel is sufficient for PromoteAlloca but is elided 167201342Snyan// before codegen. Replacing this with an equivalent intrinsic which lasts until 168201342Snyan// shortly after the machine function lowering of LDS would help break the name 169201342Snyan// mapping. The other part needed is probably to amend PromoteAlloca to embed 170201342Snyan// the LDS variables it creates in the same struct created here. That avoids the 171201342Snyan// current hazard where a PromoteAlloca LDS variable might be allocated before 172201342Snyan// the kernel scope (and thus error on the address check). Given a new invariant 173201342Snyan// that no LDS variables exist outside of the structs managed here, and an 174235988Sgleb// intrinsic that lasts until after the LDS frame lowering, it should be 175201342Snyan// possible to drop the name mapping and fold equivalent memory layouts. 176201342Snyan// 177201342Snyan//===----------------------------------------------------------------------===// 178201342Snyan 179201342Snyan#include "AMDGPU.h" 180201342Snyan#include "AMDGPUTargetMachine.h" 181201342Snyan#include "Utils/AMDGPUBaseInfo.h" 182201342Snyan#include "Utils/AMDGPUMemoryUtils.h" 183201342Snyan#include "llvm/ADT/BitVector.h" 184201342Snyan#include "llvm/ADT/DenseMap.h" 185201342Snyan#include "llvm/ADT/DenseSet.h" 186201342Snyan#include "llvm/ADT/STLExtras.h" 187201342Snyan#include "llvm/ADT/SetOperations.h" 188201342Snyan#include "llvm/Analysis/CallGraph.h" 189201342Snyan#include "llvm/CodeGen/TargetPassConfig.h" 190201342Snyan#include "llvm/IR/Constants.h" 191201342Snyan#include "llvm/IR/DerivedTypes.h" 192201342Snyan#include "llvm/IR/IRBuilder.h" 193201342Snyan#include "llvm/IR/InlineAsm.h" 194201342Snyan#include "llvm/IR/Instructions.h" 195201342Snyan#include "llvm/IR/IntrinsicsAMDGPU.h" 196201342Snyan#include "llvm/IR/MDBuilder.h" 197201342Snyan#include "llvm/IR/ReplaceConstant.h" 198201342Snyan#include "llvm/InitializePasses.h" 199201342Snyan#include "llvm/Pass.h" 200201342Snyan#include "llvm/Support/CommandLine.h" 201201342Snyan#include "llvm/Support/Debug.h" 202201342Snyan#include "llvm/Support/Format.h" 203201342Snyan#include "llvm/Support/OptimizedStructLayout.h" 204201342Snyan#include "llvm/Support/raw_ostream.h" 205201342Snyan#include "llvm/Transforms/Utils/BasicBlockUtils.h" 206201342Snyan#include "llvm/Transforms/Utils/ModuleUtils.h" 207201342Snyan 208201342Snyan#include <vector> 209201342Snyan 210201342Snyan#include <cstdio> 211201342Snyan 212201342Snyan#define DEBUG_TYPE "amdgpu-lower-module-lds" 213201342Snyan 214201342Snyanusing namespace llvm; 215201342Snyan 216201342Snyannamespace { 217201342Snyan 218201342Snyancl::opt<bool> SuperAlignLDSGlobals( 219201342Snyan "amdgpu-super-align-lds-globals", 220201342Snyan cl::desc("Increase alignment of LDS if it is not on align boundary"), 221201342Snyan cl::init(true), cl::Hidden); 222201342Snyan 223201342Snyanenum class LoweringKind { module, table, kernel, hybrid }; 224201342Snyancl::opt<LoweringKind> LoweringKindLoc( 225201342Snyan "amdgpu-lower-module-lds-strategy", 226201342Snyan cl::desc("Specify lowering strategy for function LDS access:"), cl::Hidden, 227201342Snyan cl::init(LoweringKind::hybrid), 228201342Snyan cl::values( 229201342Snyan clEnumValN(LoweringKind::table, "table", "Lower via table lookup"), 230201342Snyan clEnumValN(LoweringKind::module, "module", "Lower via module struct"), 231201342Snyan clEnumValN( 232201342Snyan LoweringKind::kernel, "kernel", 233201342Snyan "Lower variables reachable from one kernel, otherwise abort"), 234201342Snyan clEnumValN(LoweringKind::hybrid, "hybrid", 235201342Snyan "Lower via mixture of above strategies"))); 236201342Snyan 237201342Snyanbool isKernelLDS(const Function *F) { 238201342Snyan // Some weirdness here. AMDGPU::isKernelCC does not call into 239201342Snyan // AMDGPU::isKernel with the calling conv, it instead calls into 240201342Snyan // isModuleEntryFunction which returns true for more calling conventions 241201342Snyan // than AMDGPU::isKernel does. There's a FIXME on AMDGPU::isKernel. 242201342Snyan // There's also a test that checks that the LDS lowering does not hit on 243201342Snyan // a graphics shader, denoted amdgpu_ps, so stay with the limited case. 244201342Snyan // Putting LDS in the name of the function to draw attention to this. 245201342Snyan return AMDGPU::isKernel(F->getCallingConv()); 246201342Snyan} 247201342Snyan 248201342Snyantemplate <typename T> std::vector<T> sortByName(std::vector<T> &&V) { 249201342Snyan llvm::sort(V.begin(), V.end(), [](const auto *L, const auto *R) { 250201342Snyan return L->getName() < R->getName(); 251201342Snyan }); 252201342Snyan return {std::move(V)}; 253201342Snyan} 254201342Snyan 255201342Snyanclass AMDGPULowerModuleLDS { 256201342Snyan const AMDGPUTargetMachine &TM; 257201342Snyan 258201342Snyan static void 259201342Snyan removeLocalVarsFromUsedLists(Module &M, 260201342Snyan const DenseSet<GlobalVariable *> &LocalVars) { 261201342Snyan // The verifier rejects used lists containing an inttoptr of a constant 262201342Snyan // so remove the variables from these lists before replaceAllUsesWith 263201342Snyan SmallPtrSet<Constant *, 8> LocalVarsSet; 264201342Snyan for (GlobalVariable *LocalVar : LocalVars) 265201342Snyan LocalVarsSet.insert(cast<Constant>(LocalVar->stripPointerCasts())); 266201342Snyan 267201342Snyan removeFromUsedLists( 268201342Snyan M, [&LocalVarsSet](Constant *C) { return LocalVarsSet.count(C); }); 269201342Snyan 270201342Snyan for (GlobalVariable *LocalVar : LocalVars) 271201342Snyan LocalVar->removeDeadConstantUsers(); 272201342Snyan } 273201342Snyan 274201342Snyan static void markUsedByKernel(Function *Func, GlobalVariable *SGV) { 275201342Snyan // The llvm.amdgcn.module.lds instance is implicitly used by all kernels 276201342Snyan // that might call a function which accesses a field within it. This is 277201342Snyan // presently approximated to 'all kernels' if there are any such functions 278201342Snyan // in the module. This implicit use is redefined as an explicit use here so 279201342Snyan // that later passes, specifically PromoteAlloca, account for the required 280201342Snyan // memory without any knowledge of this transform. 281201342Snyan 282201342Snyan // An operand bundle on llvm.donothing works because the call instruction 283201342Snyan // survives until after the last pass that needs to account for LDS. It is 284201342Snyan // better than inline asm as the latter survives until the end of codegen. A 285201342Snyan // totally robust solution would be a function with the same semantics as 286201342Snyan // llvm.donothing that takes a pointer to the instance and is lowered to a 287201342Snyan // no-op after LDS is allocated, but that is not presently necessary. 288201342Snyan 289201342Snyan // This intrinsic is eliminated shortly before instruction selection. It 290201342Snyan // does not suffice to indicate to ISel that a given global which is not 291201342Snyan // immediately used by the kernel must still be allocated by it. An 292201342Snyan // equivalent target specific intrinsic which lasts until immediately after 293201342Snyan // codegen would suffice for that, but one would still need to ensure that 294201342Snyan // the variables are allocated in the anticpated order. 295201342Snyan BasicBlock *Entry = &Func->getEntryBlock(); 296201342Snyan IRBuilder<> Builder(Entry, Entry->getFirstNonPHIIt()); 297201342Snyan 298201342Snyan Function *Decl = 299201342Snyan Intrinsic::getDeclaration(Func->getParent(), Intrinsic::donothing, {}); 300201342Snyan 301201342Snyan Value *UseInstance[1] = { 302201342Snyan Builder.CreateConstInBoundsGEP1_32(SGV->getValueType(), SGV, 0)}; 303201342Snyan 304201342Snyan Builder.CreateCall( 305201342Snyan Decl, {}, {OperandBundleDefT<Value *>("ExplicitUse", UseInstance)}); 306201342Snyan } 307201342Snyan 308201342Snyan static bool eliminateConstantExprUsesOfLDSFromAllInstructions(Module &M) { 309201342Snyan // Constants are uniqued within LLVM. A ConstantExpr referring to a LDS 310201342Snyan // global may have uses from multiple different functions as a result. 311201342Snyan // This pass specialises LDS variables with respect to the kernel that 312201342Snyan // allocates them. 313201342Snyan 314201342Snyan // This is semantically equivalent to (the unimplemented as slow): 315201342Snyan // for (auto &F : M.functions()) 316201342Snyan // for (auto &BB : F) 317201342Snyan // for (auto &I : BB) 318201342Snyan // for (Use &Op : I.operands()) 319201342Snyan // if (constantExprUsesLDS(Op)) 320201342Snyan // replaceConstantExprInFunction(I, Op); 321201342Snyan 322201342Snyan SmallVector<Constant *> LDSGlobals; 323201342Snyan for (auto &GV : M.globals()) 324201342Snyan if (AMDGPU::isLDSVariableToLower(GV)) 325201342Snyan LDSGlobals.push_back(&GV); 326201342Snyan 327201342Snyan return convertUsersOfConstantsToInstructions(LDSGlobals); 328201342Snyan } 329201342Snyan 330201342Snyanpublic: 331201342Snyan AMDGPULowerModuleLDS(const AMDGPUTargetMachine &TM_) : TM(TM_) {} 332201342Snyan 333254015Smarcel using FunctionVariableMap = DenseMap<Function *, DenseSet<GlobalVariable *>>; 334201342Snyan 335254015Smarcel using VariableFunctionMap = DenseMap<GlobalVariable *, DenseSet<Function *>>; 336254015Smarcel 337201342Snyan static void getUsesOfLDSByFunction(CallGraph const &CG, Module &M, 338201342Snyan FunctionVariableMap &kernels, 339201342Snyan FunctionVariableMap &functions) { 340201342Snyan 341201342Snyan // Get uses from the current function, excluding uses by called functions 342201342Snyan // Two output variables to avoid walking the globals list twice 343201342Snyan for (auto &GV : M.globals()) { 344201342Snyan if (!AMDGPU::isLDSVariableToLower(GV)) { 345201342Snyan continue; 346201342Snyan } 347201342Snyan 348201342Snyan if (GV.isAbsoluteSymbolRef()) { 349201342Snyan report_fatal_error( 350201342Snyan "LDS variables with absolute addresses are unimplemented."); 351201342Snyan } 352218737Snyan 353235988Sgleb for (User *V : GV.users()) { 354232784Snyan if (auto *I = dyn_cast<Instruction>(V)) { 355201342Snyan Function *F = I->getFunction(); 356201342Snyan if (isKernelLDS(F)) { 357201342Snyan kernels[F].insert(&GV); 358201342Snyan } else { 359201342Snyan functions[F].insert(&GV); 360201342Snyan } 361201342Snyan } 362201342Snyan } 363201342Snyan } 364201342Snyan } 365201342Snyan 366201342Snyan struct LDSUsesInfoTy { 367201342Snyan FunctionVariableMap direct_access; 368201342Snyan FunctionVariableMap indirect_access; 369201342Snyan }; 370201342Snyan 371201342Snyan static LDSUsesInfoTy getTransitiveUsesOfLDS(CallGraph const &CG, Module &M) { 372201342Snyan 373201342Snyan FunctionVariableMap direct_map_kernel; 374201342Snyan FunctionVariableMap direct_map_function; 375201342Snyan getUsesOfLDSByFunction(CG, M, direct_map_kernel, direct_map_function); 376201342Snyan 377201342Snyan // Collect variables that are used by functions whose address has escaped 378201342Snyan DenseSet<GlobalVariable *> VariablesReachableThroughFunctionPointer; 379201342Snyan for (Function &F : M.functions()) { 380226506Sdes if (!isKernelLDS(&F)) 381232784Snyan if (F.hasAddressTaken(nullptr, 382232784Snyan /* IgnoreCallbackUses */ false, 383232784Snyan /* IgnoreAssumeLikeCalls */ false, 384232784Snyan /* IgnoreLLVMUsed */ true, 385201342Snyan /* IgnoreArcAttachedCall */ false)) { 386201342Snyan set_union(VariablesReachableThroughFunctionPointer, 387201342Snyan direct_map_function[&F]); 388201342Snyan } 389201342Snyan } 390201342Snyan 391201342Snyan auto functionMakesUnknownCall = [&](const Function *F) -> bool { 392201342Snyan assert(!F->isDeclaration()); 393201342Snyan for (const CallGraphNode::CallRecord &R : *CG[F]) { 394201342Snyan if (!R.second->getFunction()) { 395201342Snyan return true; 396201342Snyan } 397201342Snyan } 398201342Snyan return false; 399201342Snyan }; 400201342Snyan 401232784Snyan // Work out which variables are reachable through function calls 402219225Snyan FunctionVariableMap transitive_map_function = direct_map_function; 403232784Snyan 404201342Snyan // If the function makes any unknown call, assume the worst case that it can 405219225Snyan // access all variables accessed by functions whose address escaped 406201342Snyan for (Function &F : M.functions()) { 407201342Snyan if (!F.isDeclaration() && functionMakesUnknownCall(&F)) { 408201342Snyan if (!isKernelLDS(&F)) { 409201342Snyan set_union(transitive_map_function[&F], 410201342Snyan VariablesReachableThroughFunctionPointer); 411201342Snyan } 412201342Snyan } 413201342Snyan } 414201342Snyan 415201342Snyan // Direct implementation of collecting all variables reachable from each 416201342Snyan // function 417201342Snyan for (Function &Func : M.functions()) { 418201342Snyan if (Func.isDeclaration() || isKernelLDS(&Func)) 419201342Snyan continue; 420219225Snyan 421201342Snyan DenseSet<Function *> seen; // catches cycles 422201342Snyan SmallVector<Function *, 4> wip{&Func}; 423201342Snyan 424201342Snyan while (!wip.empty()) { 425201342Snyan Function *F = wip.pop_back_val(); 426201342Snyan 427201342Snyan // Can accelerate this by referring to transitive map for functions that 428201342Snyan // have already been computed, with more care than this 429201342Snyan set_union(transitive_map_function[&Func], direct_map_function[F]); 430201342Snyan 431201342Snyan for (const CallGraphNode::CallRecord &R : *CG[F]) { 432201342Snyan Function *ith = R.second->getFunction(); 433201342Snyan if (ith) { 434201342Snyan if (!seen.contains(ith)) { 435201342Snyan seen.insert(ith); 436201342Snyan wip.push_back(ith); 437201342Snyan } 438201342Snyan } 439201342Snyan } 440201342Snyan } 441201342Snyan } 442201342Snyan 443201342Snyan // direct_map_kernel lists which variables are used by the kernel 444201342Snyan // find the variables which are used through a function call 445201342Snyan FunctionVariableMap indirect_map_kernel; 446201342Snyan 447201342Snyan for (Function &Func : M.functions()) { 448235988Sgleb if (Func.isDeclaration() || !isKernelLDS(&Func)) 449219960Snyan continue; 450218737Snyan 451201342Snyan for (const CallGraphNode::CallRecord &R : *CG[&Func]) { 452201342Snyan Function *ith = R.second->getFunction(); 453201342Snyan if (ith) { 454201342Snyan set_union(indirect_map_kernel[&Func], transitive_map_function[ith]); 455201342Snyan } else { 456201342Snyan set_union(indirect_map_kernel[&Func], 457201342Snyan VariablesReachableThroughFunctionPointer); 458201342Snyan } 459219960Snyan } 460219960Snyan } 461201342Snyan 462201342Snyan return {std::move(direct_map_kernel), std::move(indirect_map_kernel)}; 463201342Snyan } 464201342Snyan 465201342Snyan struct LDSVariableReplacement { 466201342Snyan GlobalVariable *SGV = nullptr; 467201342Snyan DenseMap<GlobalVariable *, Constant *> LDSVarsToConstantGEP; 468201342Snyan }; 469219960Snyan 470201342Snyan // remap from lds global to a constantexpr gep to where it has been moved to 471201342Snyan // for each kernel 472201342Snyan // an array with an element for each kernel containing where the corresponding 473201342Snyan // variable was remapped to 474201342Snyan 475201342Snyan static Constant *getAddressesOfVariablesInKernel( 476201342Snyan LLVMContext &Ctx, ArrayRef<GlobalVariable *> Variables, 477201342Snyan const DenseMap<GlobalVariable *, Constant *> &LDSVarsToConstantGEP) { 478201342Snyan // Create a ConstantArray containing the address of each Variable within the 479201342Snyan // kernel corresponding to LDSVarsToConstantGEP, or poison if that kernel 480201342Snyan // does not allocate it 481201342Snyan // TODO: Drop the ptrtoint conversion 482201342Snyan 483201342Snyan Type *I32 = Type::getInt32Ty(Ctx); 484201342Snyan 485201342Snyan ArrayType *KernelOffsetsType = ArrayType::get(I32, Variables.size()); 486201342Snyan 487201342Snyan SmallVector<Constant *> Elements; 488201342Snyan for (size_t i = 0; i < Variables.size(); i++) { 489201342Snyan GlobalVariable *GV = Variables[i]; 490201342Snyan auto ConstantGepIt = LDSVarsToConstantGEP.find(GV); 491214257Snyan if (ConstantGepIt != LDSVarsToConstantGEP.end()) { 492201342Snyan auto elt = ConstantExpr::getPtrToInt(ConstantGepIt->second, I32); 493201342Snyan Elements.push_back(elt); 494201342Snyan } else { 495201342Snyan Elements.push_back(PoisonValue::get(I32)); 496201342Snyan } 497201342Snyan } 498201342Snyan return ConstantArray::get(KernelOffsetsType, Elements); 499201342Snyan } 500218842Snyan 501219960Snyan static GlobalVariable *buildLookupTable( 502219960Snyan Module &M, ArrayRef<GlobalVariable *> Variables, 503219960Snyan ArrayRef<Function *> kernels, 504201342Snyan DenseMap<Function *, LDSVariableReplacement> &KernelToReplacement) { 505219960Snyan if (Variables.empty()) { 506201342Snyan return nullptr; 507201342Snyan } 508201342Snyan LLVMContext &Ctx = M.getContext(); 509201342Snyan 510201342Snyan const size_t NumberVariables = Variables.size(); 511201342Snyan const size_t NumberKernels = kernels.size(); 512201342Snyan 513201342Snyan ArrayType *KernelOffsetsType = 514201342Snyan ArrayType::get(Type::getInt32Ty(Ctx), NumberVariables); 515201342Snyan 516201342Snyan ArrayType *AllKernelsOffsetsType = 517201342Snyan ArrayType::get(KernelOffsetsType, NumberKernels); 518201342Snyan 519201342Snyan Constant *Missing = PoisonValue::get(KernelOffsetsType); 520201342Snyan std::vector<Constant *> overallConstantExprElts(NumberKernels); 521201342Snyan for (size_t i = 0; i < NumberKernels; i++) { 522201342Snyan auto Replacement = KernelToReplacement.find(kernels[i]); 523201342Snyan overallConstantExprElts[i] = 524201342Snyan (Replacement == KernelToReplacement.end()) 525201342Snyan ? Missing 526201342Snyan : getAddressesOfVariablesInKernel( 527201342Snyan Ctx, Variables, Replacement->second.LDSVarsToConstantGEP); 528201342Snyan } 529201342Snyan 530201342Snyan Constant *init = 531201342Snyan ConstantArray::get(AllKernelsOffsetsType, overallConstantExprElts); 532201342Snyan 533201342Snyan return new GlobalVariable( 534201342Snyan M, AllKernelsOffsetsType, true, GlobalValue::InternalLinkage, init, 535201342Snyan "llvm.amdgcn.lds.offset.table", nullptr, GlobalValue::NotThreadLocal, 536201342Snyan AMDGPUAS::CONSTANT_ADDRESS); 537201342Snyan } 538201342Snyan 539201342Snyan void replaceUseWithTableLookup(Module &M, IRBuilder<> &Builder, 540201342Snyan GlobalVariable *LookupTable, 541201342Snyan GlobalVariable *GV, Use &U, 542201342Snyan Value *OptionalIndex) { 543201342Snyan // Table is a constant array of the same length as OrderedKernels 544201342Snyan LLVMContext &Ctx = M.getContext(); 545201342Snyan Type *I32 = Type::getInt32Ty(Ctx); 546201342Snyan auto *I = cast<Instruction>(U.getUser()); 547201342Snyan 548201342Snyan Value *tableKernelIndex = getTableLookupKernelIndex(M, I->getFunction()); 549201342Snyan 550201342Snyan if (auto *Phi = dyn_cast<PHINode>(I)) { 551201342Snyan BasicBlock *BB = Phi->getIncomingBlock(U); 552201342Snyan Builder.SetInsertPoint(&(*(BB->getFirstInsertionPt()))); 553201342Snyan } else { 554201342Snyan Builder.SetInsertPoint(I); 555201342Snyan } 556201342Snyan 557242863Snyan SmallVector<Value *, 3> GEPIdx = { 558242863Snyan ConstantInt::get(I32, 0), 559242863Snyan tableKernelIndex, 560242863Snyan }; 561201342Snyan if (OptionalIndex) 562201342Snyan GEPIdx.push_back(OptionalIndex); 563201342Snyan 564201342Snyan Value *Address = Builder.CreateInBoundsGEP( 565201342Snyan LookupTable->getValueType(), LookupTable, GEPIdx, GV->getName()); 566201342Snyan 567201342Snyan Value *loaded = Builder.CreateLoad(I32, Address); 568201342Snyan 569201342Snyan Value *replacement = 570201342Snyan Builder.CreateIntToPtr(loaded, GV->getType(), GV->getName()); 571201342Snyan 572201342Snyan U.set(replacement); 573201342Snyan } 574201342Snyan 575201342Snyan void replaceUsesInInstructionsWithTableLookup( 576201342Snyan Module &M, ArrayRef<GlobalVariable *> ModuleScopeVariables, 577201342Snyan GlobalVariable *LookupTable) { 578201342Snyan 579201342Snyan LLVMContext &Ctx = M.getContext(); 580201342Snyan IRBuilder<> Builder(Ctx); 581201342Snyan Type *I32 = Type::getInt32Ty(Ctx); 582201342Snyan 583201342Snyan for (size_t Index = 0; Index < ModuleScopeVariables.size(); Index++) { 584201342Snyan auto *GV = ModuleScopeVariables[Index]; 585201342Snyan 586254015Smarcel for (Use &U : make_early_inc_range(GV->uses())) { 587201342Snyan auto *I = dyn_cast<Instruction>(U.getUser()); 588201342Snyan if (!I) 589201342Snyan continue; 590201342Snyan 591201342Snyan replaceUseWithTableLookup(M, Builder, LookupTable, GV, U, 592201342Snyan ConstantInt::get(I32, Index)); 593201342Snyan } 594201342Snyan } 595201342Snyan } 596201342Snyan 597201342Snyan static DenseSet<Function *> kernelsThatIndirectlyAccessAnyOfPassedVariables( 598201342Snyan Module &M, LDSUsesInfoTy &LDSUsesInfo, 599201342Snyan DenseSet<GlobalVariable *> const &VariableSet) { 600201342Snyan 601201342Snyan DenseSet<Function *> KernelSet; 602232784Snyan 603232784Snyan if (VariableSet.empty()) 604232784Snyan return KernelSet; 605232784Snyan 606232784Snyan for (Function &Func : M.functions()) { 607232784Snyan if (Func.isDeclaration() || !isKernelLDS(&Func)) 608201342Snyan continue; 609201342Snyan for (GlobalVariable *GV : LDSUsesInfo.indirect_access[&Func]) { 610201342Snyan if (VariableSet.contains(GV)) { 611201342Snyan KernelSet.insert(&Func); 612201342Snyan break; 613201342Snyan } 614201342Snyan } 615201342Snyan } 616201342Snyan 617201342Snyan return KernelSet; 618201342Snyan } 619201342Snyan 620239063Snyan static GlobalVariable * 621239063Snyan chooseBestVariableForModuleStrategy(const DataLayout &DL, 622201342Snyan VariableFunctionMap &LDSVars) { 623201342Snyan // Find the global variable with the most indirect uses from kernels 624201342Snyan 625201342Snyan struct CandidateTy { 626201342Snyan GlobalVariable *GV = nullptr; 627201342Snyan size_t UserCount = 0; 628201342Snyan size_t Size = 0; 629254015Smarcel 630201342Snyan CandidateTy() = default; 631254015Smarcel 632201342Snyan CandidateTy(GlobalVariable *GV, uint64_t UserCount, uint64_t AllocSize) 633201342Snyan : GV(GV), UserCount(UserCount), Size(AllocSize) {} 634254015Smarcel 635201342Snyan bool operator<(const CandidateTy &Other) const { 636201342Snyan // Fewer users makes module scope variable less attractive 637201342Snyan if (UserCount < Other.UserCount) { 638201342Snyan return true; 639201342Snyan } 640201342Snyan if (UserCount > Other.UserCount) { 641201342Snyan return false; 642201342Snyan } 643201342Snyan 644201342Snyan // Bigger makes module scope variable less attractive 645201342Snyan if (Size < Other.Size) { 646201342Snyan return false; 647201342Snyan } 648201342Snyan 649201342Snyan if (Size > Other.Size) { 650201342Snyan return true; 651201342Snyan } 652201342Snyan 653201342Snyan // Arbitrary but consistent 654201342Snyan return GV->getName() < Other.GV->getName(); 655201342Snyan } 656201342Snyan }; 657201342Snyan 658201342Snyan CandidateTy MostUsed; 659201342Snyan 660201342Snyan for (auto &K : LDSVars) { 661201342Snyan GlobalVariable *GV = K.first; 662201342Snyan if (K.second.size() <= 1) { 663201342Snyan // A variable reachable by only one kernel is best lowered with kernel 664201342Snyan // strategy 665201342Snyan continue; 666201342Snyan } 667201342Snyan CandidateTy Candidate( 668201342Snyan GV, K.second.size(), 669201342Snyan DL.getTypeAllocSize(GV->getValueType()).getFixedValue()); 670201342Snyan if (MostUsed < Candidate) 671201342Snyan MostUsed = Candidate; 672201342Snyan } 673201342Snyan 674201342Snyan return MostUsed.GV; 675201342Snyan } 676201342Snyan 677201342Snyan static void recordLDSAbsoluteAddress(Module *M, GlobalVariable *GV, 678201342Snyan uint32_t Address) { 679201342Snyan // Write the specified address into metadata where it can be retrieved by 680219960Snyan // the assembler. Format is a half open range, [Address Address+1) 681201342Snyan LLVMContext &Ctx = M->getContext(); 682201342Snyan auto *IntTy = 683201342Snyan M->getDataLayout().getIntPtrType(Ctx, AMDGPUAS::LOCAL_ADDRESS); 684201342Snyan auto *MinC = ConstantAsMetadata::get(ConstantInt::get(IntTy, Address)); 685201342Snyan auto *MaxC = ConstantAsMetadata::get(ConstantInt::get(IntTy, Address + 1)); 686201342Snyan GV->setMetadata(LLVMContext::MD_absolute_symbol, 687201342Snyan MDNode::get(Ctx, {MinC, MaxC})); 688201342Snyan } 689201342Snyan 690201342Snyan DenseMap<Function *, Value *> tableKernelIndexCache; 691201342Snyan Value *getTableLookupKernelIndex(Module &M, Function *F) { 692201342Snyan // Accesses from a function use the amdgcn_lds_kernel_id intrinsic which 693201342Snyan // lowers to a read from a live in register. Emit it once in the entry 694201342Snyan // block to spare deduplicating it later. 695201342Snyan auto [It, Inserted] = tableKernelIndexCache.try_emplace(F); 696201342Snyan if (Inserted) { 697201342Snyan Function *Decl = 698201342Snyan Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_lds_kernel_id, {}); 699201342Snyan 700201342Snyan auto InsertAt = F->getEntryBlock().getFirstNonPHIOrDbgOrAlloca(); 701201342Snyan IRBuilder<> Builder(&*InsertAt); 702201342Snyan 703201342Snyan It->second = Builder.CreateCall(Decl, {}); 704201342Snyan } 705201342Snyan 706201342Snyan return It->second; 707201342Snyan } 708201342Snyan 709201342Snyan static std::vector<Function *> assignLDSKernelIDToEachKernel( 710201342Snyan Module *M, DenseSet<Function *> const &KernelsThatAllocateTableLDS, 711201342Snyan DenseSet<Function *> const &KernelsThatIndirectlyAllocateDynamicLDS) { 712201342Snyan // Associate kernels in the set with an arbirary but reproducible order and 713201342Snyan // annotate them with that order in metadata. This metadata is recognised by 714201342Snyan // the backend and lowered to a SGPR which can be read from using 715201342Snyan // amdgcn_lds_kernel_id. 716201342Snyan 717201342Snyan std::vector<Function *> OrderedKernels; 718201342Snyan if (!KernelsThatAllocateTableLDS.empty() || 719201342Snyan !KernelsThatIndirectlyAllocateDynamicLDS.empty()) { 720201342Snyan 721201342Snyan for (Function &Func : M->functions()) { 722201342Snyan if (Func.isDeclaration()) 723201342Snyan continue; 724201342Snyan if (!isKernelLDS(&Func)) 725201342Snyan continue; 726201342Snyan 727201342Snyan if (KernelsThatAllocateTableLDS.contains(&Func) || 728201342Snyan KernelsThatIndirectlyAllocateDynamicLDS.contains(&Func)) { 729201342Snyan assert(Func.hasName()); // else fatal error earlier 730201342Snyan OrderedKernels.push_back(&Func); 731201342Snyan } 732201342Snyan } 733201342Snyan 734201342Snyan // Put them in an arbitrary but reproducible order 735201342Snyan OrderedKernels = sortByName(std::move(OrderedKernels)); 736201342Snyan 737201342Snyan // Annotate the kernels with their order in this vector 738201342Snyan LLVMContext &Ctx = M->getContext(); 739201342Snyan IRBuilder<> Builder(Ctx); 740201342Snyan 741201342Snyan if (OrderedKernels.size() > UINT32_MAX) { 742201342Snyan // 32 bit keeps it in one SGPR. > 2**32 kernels won't fit on the GPU 743201342Snyan report_fatal_error("Unimplemented LDS lowering for > 2**32 kernels"); 744201342Snyan } 745201342Snyan 746201342Snyan for (size_t i = 0; i < OrderedKernels.size(); i++) { 747201342Snyan Metadata *AttrMDArgs[1] = { 748201342Snyan ConstantAsMetadata::get(Builder.getInt32(i)), 749201342Snyan }; 750201342Snyan OrderedKernels[i]->setMetadata("llvm.amdgcn.lds.kernel.id", 751201342Snyan MDNode::get(Ctx, AttrMDArgs)); 752201342Snyan } 753201342Snyan } 754201342Snyan return OrderedKernels; 755201342Snyan } 756201342Snyan 757201342Snyan static void partitionVariablesIntoIndirectStrategies( 758201342Snyan Module &M, LDSUsesInfoTy const &LDSUsesInfo, 759201342Snyan VariableFunctionMap &LDSToKernelsThatNeedToAccessItIndirectly, 760201342Snyan DenseSet<GlobalVariable *> &ModuleScopeVariables, 761201342Snyan DenseSet<GlobalVariable *> &TableLookupVariables, 762201342Snyan DenseSet<GlobalVariable *> &KernelAccessVariables, 763201342Snyan DenseSet<GlobalVariable *> &DynamicVariables) { 764201342Snyan 765201342Snyan GlobalVariable *HybridModuleRoot = 766201342Snyan LoweringKindLoc != LoweringKind::hybrid 767201342Snyan ? nullptr 768201342Snyan : chooseBestVariableForModuleStrategy( 769201342Snyan M.getDataLayout(), LDSToKernelsThatNeedToAccessItIndirectly); 770201342Snyan 771201342Snyan DenseSet<Function *> const EmptySet; 772201342Snyan DenseSet<Function *> const &HybridModuleRootKernels = 773201342Snyan HybridModuleRoot 774201342Snyan ? LDSToKernelsThatNeedToAccessItIndirectly[HybridModuleRoot] 775201342Snyan : EmptySet; 776201342Snyan 777201342Snyan for (auto &K : LDSToKernelsThatNeedToAccessItIndirectly) { 778201342Snyan // Each iteration of this loop assigns exactly one global variable to 779201342Snyan // exactly one of the implementation strategies. 780201342Snyan 781201342Snyan GlobalVariable *GV = K.first; 782201342Snyan assert(AMDGPU::isLDSVariableToLower(*GV)); 783201342Snyan assert(K.second.size() != 0); 784201342Snyan 785201342Snyan if (AMDGPU::isDynamicLDS(*GV)) { 786201342Snyan DynamicVariables.insert(GV); 787201342Snyan continue; 788201342Snyan } 789201342Snyan 790201342Snyan switch (LoweringKindLoc) { 791220685Snyan case LoweringKind::module: 792220685Snyan ModuleScopeVariables.insert(GV); 793220685Snyan break; 794220685Snyan 795220685Snyan case LoweringKind::table: 796220685Snyan TableLookupVariables.insert(GV); 797220685Snyan break; 798220685Snyan 799220685Snyan case LoweringKind::kernel: 800220685Snyan if (K.second.size() == 1) { 801220685Snyan KernelAccessVariables.insert(GV); 802220685Snyan } else { 803201342Snyan report_fatal_error( 804201342Snyan "cannot lower LDS '" + GV->getName() + 805201342Snyan "' to kernel access as it is reachable from multiple kernels"); 806201342Snyan } 807201342Snyan break; 808201342Snyan 809201342Snyan case LoweringKind::hybrid: { 810201342Snyan if (GV == HybridModuleRoot) { 811201342Snyan assert(K.second.size() != 1); 812201342Snyan ModuleScopeVariables.insert(GV); 813201342Snyan } else if (K.second.size() == 1) { 814201342Snyan KernelAccessVariables.insert(GV); 815201342Snyan } else if (set_is_subset(K.second, HybridModuleRootKernels)) { 816 ModuleScopeVariables.insert(GV); 817 } else { 818 TableLookupVariables.insert(GV); 819 } 820 break; 821 } 822 } 823 } 824 825 // All LDS variables accessed indirectly have now been partitioned into 826 // the distinct lowering strategies. 827 assert(ModuleScopeVariables.size() + TableLookupVariables.size() + 828 KernelAccessVariables.size() + DynamicVariables.size() == 829 LDSToKernelsThatNeedToAccessItIndirectly.size()); 830 } 831 832 static GlobalVariable *lowerModuleScopeStructVariables( 833 Module &M, DenseSet<GlobalVariable *> const &ModuleScopeVariables, 834 DenseSet<Function *> const &KernelsThatAllocateModuleLDS) { 835 // Create a struct to hold the ModuleScopeVariables 836 // Replace all uses of those variables from non-kernel functions with the 837 // new struct instance Replace only the uses from kernel functions that will 838 // allocate this instance. That is a space optimisation - kernels that use a 839 // subset of the module scope struct and do not need to allocate it for 840 // indirect calls will only allocate the subset they use (they do so as part 841 // of the per-kernel lowering). 842 if (ModuleScopeVariables.empty()) { 843 return nullptr; 844 } 845 846 LLVMContext &Ctx = M.getContext(); 847 848 LDSVariableReplacement ModuleScopeReplacement = 849 createLDSVariableReplacement(M, "llvm.amdgcn.module.lds", 850 ModuleScopeVariables); 851 852 appendToCompilerUsed(M, {static_cast<GlobalValue *>( 853 ConstantExpr::getPointerBitCastOrAddrSpaceCast( 854 cast<Constant>(ModuleScopeReplacement.SGV), 855 PointerType::getUnqual(Ctx)))}); 856 857 // module.lds will be allocated at zero in any kernel that allocates it 858 recordLDSAbsoluteAddress(&M, ModuleScopeReplacement.SGV, 0); 859 860 // historic 861 removeLocalVarsFromUsedLists(M, ModuleScopeVariables); 862 863 // Replace all uses of module scope variable from non-kernel functions 864 replaceLDSVariablesWithStruct( 865 M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) { 866 Instruction *I = dyn_cast<Instruction>(U.getUser()); 867 if (!I) { 868 return false; 869 } 870 Function *F = I->getFunction(); 871 return !isKernelLDS(F); 872 }); 873 874 // Replace uses of module scope variable from kernel functions that 875 // allocate the module scope variable, otherwise leave them unchanged 876 // Record on each kernel whether the module scope global is used by it 877 878 for (Function &Func : M.functions()) { 879 if (Func.isDeclaration() || !isKernelLDS(&Func)) 880 continue; 881 882 if (KernelsThatAllocateModuleLDS.contains(&Func)) { 883 replaceLDSVariablesWithStruct( 884 M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) { 885 Instruction *I = dyn_cast<Instruction>(U.getUser()); 886 if (!I) { 887 return false; 888 } 889 Function *F = I->getFunction(); 890 return F == &Func; 891 }); 892 893 markUsedByKernel(&Func, ModuleScopeReplacement.SGV); 894 } 895 } 896 897 return ModuleScopeReplacement.SGV; 898 } 899 900 static DenseMap<Function *, LDSVariableReplacement> 901 lowerKernelScopeStructVariables( 902 Module &M, LDSUsesInfoTy &LDSUsesInfo, 903 DenseSet<GlobalVariable *> const &ModuleScopeVariables, 904 DenseSet<Function *> const &KernelsThatAllocateModuleLDS, 905 GlobalVariable *MaybeModuleScopeStruct) { 906 907 // Create a struct for each kernel for the non-module-scope variables. 908 909 DenseMap<Function *, LDSVariableReplacement> KernelToReplacement; 910 for (Function &Func : M.functions()) { 911 if (Func.isDeclaration() || !isKernelLDS(&Func)) 912 continue; 913 914 DenseSet<GlobalVariable *> KernelUsedVariables; 915 // Allocating variables that are used directly in this struct to get 916 // alignment aware allocation and predictable frame size. 917 for (auto &v : LDSUsesInfo.direct_access[&Func]) { 918 if (!AMDGPU::isDynamicLDS(*v)) { 919 KernelUsedVariables.insert(v); 920 } 921 } 922 923 // Allocating variables that are accessed indirectly so that a lookup of 924 // this struct instance can find them from nested functions. 925 for (auto &v : LDSUsesInfo.indirect_access[&Func]) { 926 if (!AMDGPU::isDynamicLDS(*v)) { 927 KernelUsedVariables.insert(v); 928 } 929 } 930 931 // Variables allocated in module lds must all resolve to that struct, 932 // not to the per-kernel instance. 933 if (KernelsThatAllocateModuleLDS.contains(&Func)) { 934 for (GlobalVariable *v : ModuleScopeVariables) { 935 KernelUsedVariables.erase(v); 936 } 937 } 938 939 if (KernelUsedVariables.empty()) { 940 // Either used no LDS, or the LDS it used was all in the module struct 941 // or dynamically sized 942 continue; 943 } 944 945 // The association between kernel function and LDS struct is done by 946 // symbol name, which only works if the function in question has a 947 // name This is not expected to be a problem in practice as kernels 948 // are called by name making anonymous ones (which are named by the 949 // backend) difficult to use. This does mean that llvm test cases need 950 // to name the kernels. 951 if (!Func.hasName()) { 952 report_fatal_error("Anonymous kernels cannot use LDS variables"); 953 } 954 955 std::string VarName = 956 (Twine("llvm.amdgcn.kernel.") + Func.getName() + ".lds").str(); 957 958 auto Replacement = 959 createLDSVariableReplacement(M, VarName, KernelUsedVariables); 960 961 // If any indirect uses, create a direct use to ensure allocation 962 // TODO: Simpler to unconditionally mark used but that regresses 963 // codegen in test/CodeGen/AMDGPU/noclobber-barrier.ll 964 auto Accesses = LDSUsesInfo.indirect_access.find(&Func); 965 if ((Accesses != LDSUsesInfo.indirect_access.end()) && 966 !Accesses->second.empty()) 967 markUsedByKernel(&Func, Replacement.SGV); 968 969 // remove preserves existing codegen 970 removeLocalVarsFromUsedLists(M, KernelUsedVariables); 971 KernelToReplacement[&Func] = Replacement; 972 973 // Rewrite uses within kernel to the new struct 974 replaceLDSVariablesWithStruct( 975 M, KernelUsedVariables, Replacement, [&Func](Use &U) { 976 Instruction *I = dyn_cast<Instruction>(U.getUser()); 977 return I && I->getFunction() == &Func; 978 }); 979 } 980 return KernelToReplacement; 981 } 982 983 static GlobalVariable * 984 buildRepresentativeDynamicLDSInstance(Module &M, LDSUsesInfoTy &LDSUsesInfo, 985 Function *func) { 986 // Create a dynamic lds variable with a name associated with the passed 987 // function that has the maximum alignment of any dynamic lds variable 988 // reachable from this kernel. Dynamic LDS is allocated after the static LDS 989 // allocation, possibly after alignment padding. The representative variable 990 // created here has the maximum alignment of any other dynamic variable 991 // reachable by that kernel. All dynamic LDS variables are allocated at the 992 // same address in each kernel in order to provide the documented aliasing 993 // semantics. Setting the alignment here allows this IR pass to accurately 994 // predict the exact constant at which it will be allocated. 995 996 assert(isKernelLDS(func)); 997 998 LLVMContext &Ctx = M.getContext(); 999 const DataLayout &DL = M.getDataLayout(); 1000 Align MaxDynamicAlignment(1); 1001 1002 auto UpdateMaxAlignment = [&MaxDynamicAlignment, &DL](GlobalVariable *GV) { 1003 if (AMDGPU::isDynamicLDS(*GV)) { 1004 MaxDynamicAlignment = 1005 std::max(MaxDynamicAlignment, AMDGPU::getAlign(DL, GV)); 1006 } 1007 }; 1008 1009 for (GlobalVariable *GV : LDSUsesInfo.indirect_access[func]) { 1010 UpdateMaxAlignment(GV); 1011 } 1012 1013 for (GlobalVariable *GV : LDSUsesInfo.direct_access[func]) { 1014 UpdateMaxAlignment(GV); 1015 } 1016 1017 assert(func->hasName()); // Checked by caller 1018 auto emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0); 1019 GlobalVariable *N = new GlobalVariable( 1020 M, emptyCharArray, false, GlobalValue::ExternalLinkage, nullptr, 1021 Twine("llvm.amdgcn." + func->getName() + ".dynlds"), nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS, 1022 false); 1023 N->setAlignment(MaxDynamicAlignment); 1024 1025 assert(AMDGPU::isDynamicLDS(*N)); 1026 return N; 1027 } 1028 1029 /// Strip "amdgpu-no-lds-kernel-id" from any functions where we may have 1030 /// introduced its use. If AMDGPUAttributor ran prior to the pass, we inferred 1031 /// the lack of llvm.amdgcn.lds.kernel.id calls. 1032 void removeNoLdsKernelIdFromReachable(CallGraph &CG, Function *KernelRoot) { 1033 KernelRoot->removeFnAttr("amdgpu-no-lds-kernel-id"); 1034 1035 SmallVector<Function *> Tmp({CG[KernelRoot]->getFunction()}); 1036 if (!Tmp.back()) 1037 return; 1038 1039 SmallPtrSet<Function *, 8> Visited; 1040 bool SeenUnknownCall = false; 1041 1042 do { 1043 Function *F = Tmp.pop_back_val(); 1044 1045 for (auto &N : *CG[F]) { 1046 if (!N.second) 1047 continue; 1048 1049 Function *Callee = N.second->getFunction(); 1050 if (!Callee) { 1051 if (!SeenUnknownCall) { 1052 SeenUnknownCall = true; 1053 1054 // If we see any indirect calls, assume nothing about potential 1055 // targets. 1056 // TODO: This could be refined to possible LDS global users. 1057 for (auto &N : *CG.getExternalCallingNode()) { 1058 Function *PotentialCallee = N.second->getFunction(); 1059 if (!isKernelLDS(PotentialCallee)) 1060 PotentialCallee->removeFnAttr("amdgpu-no-lds-kernel-id"); 1061 } 1062 1063 continue; 1064 } 1065 } 1066 1067 Callee->removeFnAttr("amdgpu-no-lds-kernel-id"); 1068 if (Visited.insert(Callee).second) 1069 Tmp.push_back(Callee); 1070 } 1071 } while (!Tmp.empty()); 1072 } 1073 1074 DenseMap<Function *, GlobalVariable *> lowerDynamicLDSVariables( 1075 Module &M, LDSUsesInfoTy &LDSUsesInfo, 1076 DenseSet<Function *> const &KernelsThatIndirectlyAllocateDynamicLDS, 1077 DenseSet<GlobalVariable *> const &DynamicVariables, 1078 std::vector<Function *> const &OrderedKernels) { 1079 DenseMap<Function *, GlobalVariable *> KernelToCreatedDynamicLDS; 1080 if (!KernelsThatIndirectlyAllocateDynamicLDS.empty()) { 1081 LLVMContext &Ctx = M.getContext(); 1082 IRBuilder<> Builder(Ctx); 1083 Type *I32 = Type::getInt32Ty(Ctx); 1084 1085 std::vector<Constant *> newDynamicLDS; 1086 1087 // Table is built in the same order as OrderedKernels 1088 for (auto &func : OrderedKernels) { 1089 1090 if (KernelsThatIndirectlyAllocateDynamicLDS.contains(func)) { 1091 assert(isKernelLDS(func)); 1092 if (!func->hasName()) { 1093 report_fatal_error("Anonymous kernels cannot use LDS variables"); 1094 } 1095 1096 GlobalVariable *N = 1097 buildRepresentativeDynamicLDSInstance(M, LDSUsesInfo, func); 1098 1099 KernelToCreatedDynamicLDS[func] = N; 1100 1101 markUsedByKernel(func, N); 1102 1103 auto emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0); 1104 auto GEP = ConstantExpr::getGetElementPtr( 1105 emptyCharArray, N, ConstantInt::get(I32, 0), true); 1106 newDynamicLDS.push_back(ConstantExpr::getPtrToInt(GEP, I32)); 1107 } else { 1108 newDynamicLDS.push_back(PoisonValue::get(I32)); 1109 } 1110 } 1111 assert(OrderedKernels.size() == newDynamicLDS.size()); 1112 1113 ArrayType *t = ArrayType::get(I32, newDynamicLDS.size()); 1114 Constant *init = ConstantArray::get(t, newDynamicLDS); 1115 GlobalVariable *table = new GlobalVariable( 1116 M, t, true, GlobalValue::InternalLinkage, init, 1117 "llvm.amdgcn.dynlds.offset.table", nullptr, 1118 GlobalValue::NotThreadLocal, AMDGPUAS::CONSTANT_ADDRESS); 1119 1120 for (GlobalVariable *GV : DynamicVariables) { 1121 for (Use &U : make_early_inc_range(GV->uses())) { 1122 auto *I = dyn_cast<Instruction>(U.getUser()); 1123 if (!I) 1124 continue; 1125 if (isKernelLDS(I->getFunction())) 1126 continue; 1127 1128 replaceUseWithTableLookup(M, Builder, table, GV, U, nullptr); 1129 } 1130 } 1131 } 1132 return KernelToCreatedDynamicLDS; 1133 } 1134 1135 bool runOnModule(Module &M) { 1136 CallGraph CG = CallGraph(M); 1137 bool Changed = superAlignLDSGlobals(M); 1138 1139 Changed |= eliminateConstantExprUsesOfLDSFromAllInstructions(M); 1140 1141 Changed = true; // todo: narrow this down 1142 1143 // For each kernel, what variables does it access directly or through 1144 // callees 1145 LDSUsesInfoTy LDSUsesInfo = getTransitiveUsesOfLDS(CG, M); 1146 1147 // For each variable accessed through callees, which kernels access it 1148 VariableFunctionMap LDSToKernelsThatNeedToAccessItIndirectly; 1149 for (auto &K : LDSUsesInfo.indirect_access) { 1150 Function *F = K.first; 1151 assert(isKernelLDS(F)); 1152 for (GlobalVariable *GV : K.second) { 1153 LDSToKernelsThatNeedToAccessItIndirectly[GV].insert(F); 1154 } 1155 } 1156 1157 // Partition variables accessed indirectly into the different strategies 1158 DenseSet<GlobalVariable *> ModuleScopeVariables; 1159 DenseSet<GlobalVariable *> TableLookupVariables; 1160 DenseSet<GlobalVariable *> KernelAccessVariables; 1161 DenseSet<GlobalVariable *> DynamicVariables; 1162 partitionVariablesIntoIndirectStrategies( 1163 M, LDSUsesInfo, LDSToKernelsThatNeedToAccessItIndirectly, 1164 ModuleScopeVariables, TableLookupVariables, KernelAccessVariables, 1165 DynamicVariables); 1166 1167 // If the kernel accesses a variable that is going to be stored in the 1168 // module instance through a call then that kernel needs to allocate the 1169 // module instance 1170 const DenseSet<Function *> KernelsThatAllocateModuleLDS = 1171 kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo, 1172 ModuleScopeVariables); 1173 const DenseSet<Function *> KernelsThatAllocateTableLDS = 1174 kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo, 1175 TableLookupVariables); 1176 1177 const DenseSet<Function *> KernelsThatIndirectlyAllocateDynamicLDS = 1178 kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo, 1179 DynamicVariables); 1180 1181 GlobalVariable *MaybeModuleScopeStruct = lowerModuleScopeStructVariables( 1182 M, ModuleScopeVariables, KernelsThatAllocateModuleLDS); 1183 1184 DenseMap<Function *, LDSVariableReplacement> KernelToReplacement = 1185 lowerKernelScopeStructVariables(M, LDSUsesInfo, ModuleScopeVariables, 1186 KernelsThatAllocateModuleLDS, 1187 MaybeModuleScopeStruct); 1188 1189 // Lower zero cost accesses to the kernel instances just created 1190 for (auto &GV : KernelAccessVariables) { 1191 auto &funcs = LDSToKernelsThatNeedToAccessItIndirectly[GV]; 1192 assert(funcs.size() == 1); // Only one kernel can access it 1193 LDSVariableReplacement Replacement = 1194 KernelToReplacement[*(funcs.begin())]; 1195 1196 DenseSet<GlobalVariable *> Vec; 1197 Vec.insert(GV); 1198 1199 replaceLDSVariablesWithStruct(M, Vec, Replacement, [](Use &U) { 1200 return isa<Instruction>(U.getUser()); 1201 }); 1202 } 1203 1204 // The ith element of this vector is kernel id i 1205 std::vector<Function *> OrderedKernels = 1206 assignLDSKernelIDToEachKernel(&M, KernelsThatAllocateTableLDS, 1207 KernelsThatIndirectlyAllocateDynamicLDS); 1208 1209 if (!KernelsThatAllocateTableLDS.empty()) { 1210 LLVMContext &Ctx = M.getContext(); 1211 IRBuilder<> Builder(Ctx); 1212 1213 // The order must be consistent between lookup table and accesses to 1214 // lookup table 1215 auto TableLookupVariablesOrdered = 1216 sortByName(std::vector<GlobalVariable *>(TableLookupVariables.begin(), 1217 TableLookupVariables.end())); 1218 1219 GlobalVariable *LookupTable = buildLookupTable( 1220 M, TableLookupVariablesOrdered, OrderedKernels, KernelToReplacement); 1221 replaceUsesInInstructionsWithTableLookup(M, TableLookupVariablesOrdered, 1222 LookupTable); 1223 1224 // Strip amdgpu-no-lds-kernel-id from all functions reachable from the 1225 // kernel. We may have inferred this wasn't used prior to the pass. 1226 // 1227 // TODO: We could filter out subgraphs that do not access LDS globals. 1228 for (Function *F : KernelsThatAllocateTableLDS) 1229 removeNoLdsKernelIdFromReachable(CG, F); 1230 } 1231 1232 DenseMap<Function *, GlobalVariable *> KernelToCreatedDynamicLDS = 1233 lowerDynamicLDSVariables(M, LDSUsesInfo, 1234 KernelsThatIndirectlyAllocateDynamicLDS, 1235 DynamicVariables, OrderedKernels); 1236 1237 // All kernel frames have been allocated. Calculate and record the 1238 // addresses. 1239 { 1240 const DataLayout &DL = M.getDataLayout(); 1241 1242 for (Function &Func : M.functions()) { 1243 if (Func.isDeclaration() || !isKernelLDS(&Func)) 1244 continue; 1245 1246 // All three of these are optional. The first variable is allocated at 1247 // zero. They are allocated by AMDGPUMachineFunction as one block. 1248 // Layout: 1249 //{ 1250 // module.lds 1251 // alignment padding 1252 // kernel instance 1253 // alignment padding 1254 // dynamic lds variables 1255 //} 1256 1257 const bool AllocateModuleScopeStruct = 1258 MaybeModuleScopeStruct && 1259 KernelsThatAllocateModuleLDS.contains(&Func); 1260 1261 auto Replacement = KernelToReplacement.find(&Func); 1262 const bool AllocateKernelScopeStruct = 1263 Replacement != KernelToReplacement.end(); 1264 1265 const bool AllocateDynamicVariable = 1266 KernelToCreatedDynamicLDS.contains(&Func); 1267 1268 uint32_t Offset = 0; 1269 1270 if (AllocateModuleScopeStruct) { 1271 // Allocated at zero, recorded once on construction, not once per 1272 // kernel 1273 Offset += DL.getTypeAllocSize(MaybeModuleScopeStruct->getValueType()); 1274 } 1275 1276 if (AllocateKernelScopeStruct) { 1277 GlobalVariable *KernelStruct = Replacement->second.SGV; 1278 Offset = alignTo(Offset, AMDGPU::getAlign(DL, KernelStruct)); 1279 recordLDSAbsoluteAddress(&M, KernelStruct, Offset); 1280 Offset += DL.getTypeAllocSize(KernelStruct->getValueType()); 1281 } 1282 1283 // If there is dynamic allocation, the alignment needed is included in 1284 // the static frame size. There may be no reference to the dynamic 1285 // variable in the kernel itself, so without including it here, that 1286 // alignment padding could be missed. 1287 if (AllocateDynamicVariable) { 1288 GlobalVariable *DynamicVariable = KernelToCreatedDynamicLDS[&Func]; 1289 Offset = alignTo(Offset, AMDGPU::getAlign(DL, DynamicVariable)); 1290 recordLDSAbsoluteAddress(&M, DynamicVariable, Offset); 1291 } 1292 1293 if (Offset != 0) { 1294 (void)TM; // TODO: Account for target maximum LDS 1295 std::string Buffer; 1296 raw_string_ostream SS{Buffer}; 1297 SS << format("%u", Offset); 1298 1299 // Instead of explictly marking kernels that access dynamic variables 1300 // using special case metadata, annotate with min-lds == max-lds, i.e. 1301 // that there is no more space available for allocating more static 1302 // LDS variables. That is the right condition to prevent allocating 1303 // more variables which would collide with the addresses assigned to 1304 // dynamic variables. 1305 if (AllocateDynamicVariable) 1306 SS << format(",%u", Offset); 1307 1308 Func.addFnAttr("amdgpu-lds-size", Buffer); 1309 } 1310 } 1311 } 1312 1313 for (auto &GV : make_early_inc_range(M.globals())) 1314 if (AMDGPU::isLDSVariableToLower(GV)) { 1315 // probably want to remove from used lists 1316 GV.removeDeadConstantUsers(); 1317 if (GV.use_empty()) 1318 GV.eraseFromParent(); 1319 } 1320 1321 return Changed; 1322 } 1323 1324private: 1325 // Increase the alignment of LDS globals if necessary to maximise the chance 1326 // that we can use aligned LDS instructions to access them. 1327 static bool superAlignLDSGlobals(Module &M) { 1328 const DataLayout &DL = M.getDataLayout(); 1329 bool Changed = false; 1330 if (!SuperAlignLDSGlobals) { 1331 return Changed; 1332 } 1333 1334 for (auto &GV : M.globals()) { 1335 if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) { 1336 // Only changing alignment of LDS variables 1337 continue; 1338 } 1339 if (!GV.hasInitializer()) { 1340 // cuda/hip extern __shared__ variable, leave alignment alone 1341 continue; 1342 } 1343 1344 Align Alignment = AMDGPU::getAlign(DL, &GV); 1345 TypeSize GVSize = DL.getTypeAllocSize(GV.getValueType()); 1346 1347 if (GVSize > 8) { 1348 // We might want to use a b96 or b128 load/store 1349 Alignment = std::max(Alignment, Align(16)); 1350 } else if (GVSize > 4) { 1351 // We might want to use a b64 load/store 1352 Alignment = std::max(Alignment, Align(8)); 1353 } else if (GVSize > 2) { 1354 // We might want to use a b32 load/store 1355 Alignment = std::max(Alignment, Align(4)); 1356 } else if (GVSize > 1) { 1357 // We might want to use a b16 load/store 1358 Alignment = std::max(Alignment, Align(2)); 1359 } 1360 1361 if (Alignment != AMDGPU::getAlign(DL, &GV)) { 1362 Changed = true; 1363 GV.setAlignment(Alignment); 1364 } 1365 } 1366 return Changed; 1367 } 1368 1369 static LDSVariableReplacement createLDSVariableReplacement( 1370 Module &M, std::string VarName, 1371 DenseSet<GlobalVariable *> const &LDSVarsToTransform) { 1372 // Create a struct instance containing LDSVarsToTransform and map from those 1373 // variables to ConstantExprGEP 1374 // Variables may be introduced to meet alignment requirements. No aliasing 1375 // metadata is useful for these as they have no uses. Erased before return. 1376 1377 LLVMContext &Ctx = M.getContext(); 1378 const DataLayout &DL = M.getDataLayout(); 1379 assert(!LDSVarsToTransform.empty()); 1380 1381 SmallVector<OptimizedStructLayoutField, 8> LayoutFields; 1382 LayoutFields.reserve(LDSVarsToTransform.size()); 1383 { 1384 // The order of fields in this struct depends on the order of 1385 // varables in the argument which varies when changing how they 1386 // are identified, leading to spurious test breakage. 1387 auto Sorted = sortByName(std::vector<GlobalVariable *>( 1388 LDSVarsToTransform.begin(), LDSVarsToTransform.end())); 1389 1390 for (GlobalVariable *GV : Sorted) { 1391 OptimizedStructLayoutField F(GV, 1392 DL.getTypeAllocSize(GV->getValueType()), 1393 AMDGPU::getAlign(DL, GV)); 1394 LayoutFields.emplace_back(F); 1395 } 1396 } 1397 1398 performOptimizedStructLayout(LayoutFields); 1399 1400 std::vector<GlobalVariable *> LocalVars; 1401 BitVector IsPaddingField; 1402 LocalVars.reserve(LDSVarsToTransform.size()); // will be at least this large 1403 IsPaddingField.reserve(LDSVarsToTransform.size()); 1404 { 1405 uint64_t CurrentOffset = 0; 1406 for (size_t I = 0; I < LayoutFields.size(); I++) { 1407 GlobalVariable *FGV = static_cast<GlobalVariable *>( 1408 const_cast<void *>(LayoutFields[I].Id)); 1409 Align DataAlign = LayoutFields[I].Alignment; 1410 1411 uint64_t DataAlignV = DataAlign.value(); 1412 if (uint64_t Rem = CurrentOffset % DataAlignV) { 1413 uint64_t Padding = DataAlignV - Rem; 1414 1415 // Append an array of padding bytes to meet alignment requested 1416 // Note (o + (a - (o % a)) ) % a == 0 1417 // (offset + Padding ) % align == 0 1418 1419 Type *ATy = ArrayType::get(Type::getInt8Ty(Ctx), Padding); 1420 LocalVars.push_back(new GlobalVariable( 1421 M, ATy, false, GlobalValue::InternalLinkage, 1422 PoisonValue::get(ATy), "", nullptr, GlobalValue::NotThreadLocal, 1423 AMDGPUAS::LOCAL_ADDRESS, false)); 1424 IsPaddingField.push_back(true); 1425 CurrentOffset += Padding; 1426 } 1427 1428 LocalVars.push_back(FGV); 1429 IsPaddingField.push_back(false); 1430 CurrentOffset += LayoutFields[I].Size; 1431 } 1432 } 1433 1434 std::vector<Type *> LocalVarTypes; 1435 LocalVarTypes.reserve(LocalVars.size()); 1436 std::transform( 1437 LocalVars.cbegin(), LocalVars.cend(), std::back_inserter(LocalVarTypes), 1438 [](const GlobalVariable *V) -> Type * { return V->getValueType(); }); 1439 1440 StructType *LDSTy = StructType::create(Ctx, LocalVarTypes, VarName + ".t"); 1441 1442 Align StructAlign = AMDGPU::getAlign(DL, LocalVars[0]); 1443 1444 GlobalVariable *SGV = new GlobalVariable( 1445 M, LDSTy, false, GlobalValue::InternalLinkage, PoisonValue::get(LDSTy), 1446 VarName, nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS, 1447 false); 1448 SGV->setAlignment(StructAlign); 1449 1450 DenseMap<GlobalVariable *, Constant *> Map; 1451 Type *I32 = Type::getInt32Ty(Ctx); 1452 for (size_t I = 0; I < LocalVars.size(); I++) { 1453 GlobalVariable *GV = LocalVars[I]; 1454 Constant *GEPIdx[] = {ConstantInt::get(I32, 0), ConstantInt::get(I32, I)}; 1455 Constant *GEP = ConstantExpr::getGetElementPtr(LDSTy, SGV, GEPIdx, true); 1456 if (IsPaddingField[I]) { 1457 assert(GV->use_empty()); 1458 GV->eraseFromParent(); 1459 } else { 1460 Map[GV] = GEP; 1461 } 1462 } 1463 assert(Map.size() == LDSVarsToTransform.size()); 1464 return {SGV, std::move(Map)}; 1465 } 1466 1467 template <typename PredicateTy> 1468 static void replaceLDSVariablesWithStruct( 1469 Module &M, DenseSet<GlobalVariable *> const &LDSVarsToTransformArg, 1470 const LDSVariableReplacement &Replacement, PredicateTy Predicate) { 1471 LLVMContext &Ctx = M.getContext(); 1472 const DataLayout &DL = M.getDataLayout(); 1473 1474 // A hack... we need to insert the aliasing info in a predictable order for 1475 // lit tests. Would like to have them in a stable order already, ideally the 1476 // same order they get allocated, which might mean an ordered set container 1477 auto LDSVarsToTransform = sortByName(std::vector<GlobalVariable *>( 1478 LDSVarsToTransformArg.begin(), LDSVarsToTransformArg.end())); 1479 1480 // Create alias.scope and their lists. Each field in the new structure 1481 // does not alias with all other fields. 1482 SmallVector<MDNode *> AliasScopes; 1483 SmallVector<Metadata *> NoAliasList; 1484 const size_t NumberVars = LDSVarsToTransform.size(); 1485 if (NumberVars > 1) { 1486 MDBuilder MDB(Ctx); 1487 AliasScopes.reserve(NumberVars); 1488 MDNode *Domain = MDB.createAnonymousAliasScopeDomain(); 1489 for (size_t I = 0; I < NumberVars; I++) { 1490 MDNode *Scope = MDB.createAnonymousAliasScope(Domain); 1491 AliasScopes.push_back(Scope); 1492 } 1493 NoAliasList.append(&AliasScopes[1], AliasScopes.end()); 1494 } 1495 1496 // Replace uses of ith variable with a constantexpr to the corresponding 1497 // field of the instance that will be allocated by AMDGPUMachineFunction 1498 for (size_t I = 0; I < NumberVars; I++) { 1499 GlobalVariable *GV = LDSVarsToTransform[I]; 1500 Constant *GEP = Replacement.LDSVarsToConstantGEP.at(GV); 1501 1502 GV->replaceUsesWithIf(GEP, Predicate); 1503 1504 APInt APOff(DL.getIndexTypeSizeInBits(GEP->getType()), 0); 1505 GEP->stripAndAccumulateInBoundsConstantOffsets(DL, APOff); 1506 uint64_t Offset = APOff.getZExtValue(); 1507 1508 Align A = 1509 commonAlignment(Replacement.SGV->getAlign().valueOrOne(), Offset); 1510 1511 if (I) 1512 NoAliasList[I - 1] = AliasScopes[I - 1]; 1513 MDNode *NoAlias = 1514 NoAliasList.empty() ? nullptr : MDNode::get(Ctx, NoAliasList); 1515 MDNode *AliasScope = 1516 AliasScopes.empty() ? nullptr : MDNode::get(Ctx, {AliasScopes[I]}); 1517 1518 refineUsesAlignmentAndAA(GEP, A, DL, AliasScope, NoAlias); 1519 } 1520 } 1521 1522 static void refineUsesAlignmentAndAA(Value *Ptr, Align A, 1523 const DataLayout &DL, MDNode *AliasScope, 1524 MDNode *NoAlias, unsigned MaxDepth = 5) { 1525 if (!MaxDepth || (A == 1 && !AliasScope)) 1526 return; 1527 1528 for (User *U : Ptr->users()) { 1529 if (auto *I = dyn_cast<Instruction>(U)) { 1530 if (AliasScope && I->mayReadOrWriteMemory()) { 1531 MDNode *AS = I->getMetadata(LLVMContext::MD_alias_scope); 1532 AS = (AS ? MDNode::getMostGenericAliasScope(AS, AliasScope) 1533 : AliasScope); 1534 I->setMetadata(LLVMContext::MD_alias_scope, AS); 1535 1536 MDNode *NA = I->getMetadata(LLVMContext::MD_noalias); 1537 NA = (NA ? MDNode::intersect(NA, NoAlias) : NoAlias); 1538 I->setMetadata(LLVMContext::MD_noalias, NA); 1539 } 1540 } 1541 1542 if (auto *LI = dyn_cast<LoadInst>(U)) { 1543 LI->setAlignment(std::max(A, LI->getAlign())); 1544 continue; 1545 } 1546 if (auto *SI = dyn_cast<StoreInst>(U)) { 1547 if (SI->getPointerOperand() == Ptr) 1548 SI->setAlignment(std::max(A, SI->getAlign())); 1549 continue; 1550 } 1551 if (auto *AI = dyn_cast<AtomicRMWInst>(U)) { 1552 // None of atomicrmw operations can work on pointers, but let's 1553 // check it anyway in case it will or we will process ConstantExpr. 1554 if (AI->getPointerOperand() == Ptr) 1555 AI->setAlignment(std::max(A, AI->getAlign())); 1556 continue; 1557 } 1558 if (auto *AI = dyn_cast<AtomicCmpXchgInst>(U)) { 1559 if (AI->getPointerOperand() == Ptr) 1560 AI->setAlignment(std::max(A, AI->getAlign())); 1561 continue; 1562 } 1563 if (auto *GEP = dyn_cast<GetElementPtrInst>(U)) { 1564 unsigned BitWidth = DL.getIndexTypeSizeInBits(GEP->getType()); 1565 APInt Off(BitWidth, 0); 1566 if (GEP->getPointerOperand() == Ptr) { 1567 Align GA; 1568 if (GEP->accumulateConstantOffset(DL, Off)) 1569 GA = commonAlignment(A, Off.getLimitedValue()); 1570 refineUsesAlignmentAndAA(GEP, GA, DL, AliasScope, NoAlias, 1571 MaxDepth - 1); 1572 } 1573 continue; 1574 } 1575 if (auto *I = dyn_cast<Instruction>(U)) { 1576 if (I->getOpcode() == Instruction::BitCast || 1577 I->getOpcode() == Instruction::AddrSpaceCast) 1578 refineUsesAlignmentAndAA(I, A, DL, AliasScope, NoAlias, MaxDepth - 1); 1579 } 1580 } 1581 } 1582}; 1583 1584class AMDGPULowerModuleLDSLegacy : public ModulePass { 1585public: 1586 const AMDGPUTargetMachine *TM; 1587 static char ID; 1588 1589 AMDGPULowerModuleLDSLegacy(const AMDGPUTargetMachine *TM_ = nullptr) 1590 : ModulePass(ID), TM(TM_) { 1591 initializeAMDGPULowerModuleLDSLegacyPass(*PassRegistry::getPassRegistry()); 1592 } 1593 1594 void getAnalysisUsage(AnalysisUsage &AU) const override { 1595 if (!TM) 1596 AU.addRequired<TargetPassConfig>(); 1597 } 1598 1599 bool runOnModule(Module &M) override { 1600 if (!TM) { 1601 auto &TPC = getAnalysis<TargetPassConfig>(); 1602 TM = &TPC.getTM<AMDGPUTargetMachine>(); 1603 } 1604 1605 return AMDGPULowerModuleLDS(*TM).runOnModule(M); 1606 } 1607}; 1608 1609} // namespace 1610char AMDGPULowerModuleLDSLegacy::ID = 0; 1611 1612char &llvm::AMDGPULowerModuleLDSLegacyPassID = AMDGPULowerModuleLDSLegacy::ID; 1613 1614INITIALIZE_PASS_BEGIN(AMDGPULowerModuleLDSLegacy, DEBUG_TYPE, 1615 "Lower uses of LDS variables from non-kernel functions", 1616 false, false) 1617INITIALIZE_PASS_DEPENDENCY(TargetPassConfig) 1618INITIALIZE_PASS_END(AMDGPULowerModuleLDSLegacy, DEBUG_TYPE, 1619 "Lower uses of LDS variables from non-kernel functions", 1620 false, false) 1621 1622ModulePass * 1623llvm::createAMDGPULowerModuleLDSLegacyPass(const AMDGPUTargetMachine *TM) { 1624 return new AMDGPULowerModuleLDSLegacy(TM); 1625} 1626 1627PreservedAnalyses AMDGPULowerModuleLDSPass::run(Module &M, 1628 ModuleAnalysisManager &) { 1629 return AMDGPULowerModuleLDS(TM).runOnModule(M) ? PreservedAnalyses::none() 1630 : PreservedAnalyses::all(); 1631} 1632