1455 lines · cpp
1//===-- AMDGPULowerModuleLDSPass.cpp ------------------------------*- C++ -*-=//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// This pass eliminates local data store, LDS, uses from non-kernel functions.10// LDS is contiguous memory allocated per kernel execution.11//12// Background.13//14// The programming model is global variables, or equivalently function local15// static variables, accessible from kernels or other functions. For uses from16// kernels this is straightforward - assign an integer to the kernel for the17// memory required by all the variables combined, allocate them within that.18// For uses from functions there are performance tradeoffs to choose between.19//20// This model means the GPU runtime can specify the amount of memory allocated.21// If this is more than the kernel assumed, the excess can be made available22// using a language specific feature, which IR represents as a variable with23// no initializer. This feature is referred to here as "Dynamic LDS" and is24// lowered slightly differently to the normal case.25//26// Consequences of this GPU feature:27// - memory is limited and exceeding it halts compilation28// - a global accessed by one kernel exists independent of other kernels29// - a global exists independent of simultaneous execution of the same kernel30// - the address of the global may be different from different kernels as they31// do not alias, which permits only allocating variables they use32// - if the address is allowed to differ, functions need help to find it33//34// Uses from kernels are implemented here by grouping them in a per-kernel35// struct instance. This duplicates the variables, accurately modelling their36// aliasing properties relative to a single global representation. It also37// permits control over alignment via padding.38//39// Uses from functions are more complicated and the primary purpose of this40// IR pass. Several different lowering are chosen between to meet requirements41// to avoid allocating any LDS where it is not necessary, as that impacts42// occupancy and may fail the compilation, while not imposing overhead on a43// feature whose primary advantage over global memory is performance. The basic44// design goal is to avoid one kernel imposing overhead on another.45//46// Implementation.47//48// LDS variables with constant annotation or non-undef initializer are passed49// through unchanged for simplification or error diagnostics in later passes.50// Non-undef initializers are not yet implemented for LDS.51//52// LDS variables that are always allocated at the same address can be found53// by lookup at that address. Otherwise runtime information/cost is required.54//55// The simplest strategy possible is to group all LDS variables in a single56// struct and allocate that struct in every kernel such that the original57// variables are always at the same address. LDS is however a limited resource58// so this strategy is unusable in practice. It is not implemented here.59//60// Strategy | Precise allocation | Zero runtime cost | General purpose |61// --------+--------------------+-------------------+-----------------+62// Module | No | Yes | Yes |63// Table | Yes | No | Yes |64// Kernel | Yes | Yes | No |65// Hybrid | Yes | Partial | Yes |66//67// "Module" spends LDS memory to save cycles. "Table" spends cycles and global68// memory to save LDS. "Kernel" is as fast as kernel allocation but only works69// for variables that are known reachable from a single kernel. "Hybrid" picks70// between all three. When forced to choose between LDS and cycles we minimise71// LDS use.72 73// The "module" lowering implemented here finds LDS variables which are used by74// non-kernel functions and creates a new struct with a field for each of those75// LDS variables. Variables that are only used from kernels are excluded.76//77// The "table" lowering implemented here has three components.78// First kernels are assigned a unique integer identifier which is available in79// functions it calls through the intrinsic amdgcn_lds_kernel_id. The integer80// is passed through a specific SGPR, thus works with indirect calls.81// Second, each kernel allocates LDS variables independent of other kernels and82// writes the addresses it chose for each variable into an array in consistent83// order. If the kernel does not allocate a given variable, it writes undef to84// the corresponding array location. These arrays are written to a constant85// table in the order matching the kernel unique integer identifier.86// Third, uses from non-kernel functions are replaced with a table lookup using87// the intrinsic function to find the address of the variable.88//89// "Kernel" lowering is only applicable for variables that are unambiguously90// reachable from exactly one kernel. For those cases, accesses to the variable91// can be lowered to ConstantExpr address of a struct instance specific to that92// one kernel. This is zero cost in space and in compute. It will raise a fatal93// error on any variable that might be reachable from multiple kernels and is94// thus most easily used as part of the hybrid lowering strategy.95//96// Hybrid lowering is a mixture of the above. It uses the zero cost kernel97// lowering where it can. It lowers the variable accessed by the greatest98// number of kernels using the module strategy as that is free for the first99// variable. Any futher variables that can be lowered with the module strategy100// without incurring LDS memory overhead are. The remaining ones are lowered101// via table.102//103// Consequences104// - No heuristics or user controlled magic numbers, hybrid is the right choice105// - Kernels that don't use functions (or have had them all inlined) are not106// affected by any lowering for kernels that do.107// - Kernels that don't make indirect function calls are not affected by those108// that do.109// - Variables which are used by lots of kernels, e.g. those injected by a110// language runtime in most kernels, are expected to have no overhead111// - Implementations that instantiate templates per-kernel where those templates112// use LDS are expected to hit the "Kernel" lowering strategy113// - The runtime properties impose a cost in compiler implementation complexity114//115// Dynamic LDS implementation116// Dynamic LDS is lowered similarly to the "table" strategy above and uses the117// same intrinsic to identify which kernel is at the root of the dynamic call118// graph. This relies on the specified behaviour that all dynamic LDS variables119// alias one another, i.e. are at the same address, with respect to a given120// kernel. Therefore this pass creates new dynamic LDS variables for each kernel121// that allocates any dynamic LDS and builds a table of addresses out of those.122// The AMDGPUPromoteAlloca pass skips kernels that use dynamic LDS.123// The corresponding optimisation for "kernel" lowering where the table lookup124// is elided is not implemented.125//126//127// Implementation notes / limitations128// A single LDS global variable represents an instance per kernel that can reach129// said variables. This pass essentially specialises said variables per kernel.130// Handling ConstantExpr during the pass complicated this significantly so now131// all ConstantExpr uses of LDS variables are expanded to instructions. This132// may need amending when implementing non-undef initialisers.133//134// Lowering is split between this IR pass and the back end. This pass chooses135// where given variables should be allocated and marks them with metadata,136// MD_absolute_symbol. The backend places the variables in coincidentally the137// same location and raises a fatal error if something has gone awry. This works138// in practice because the only pass between this one and the backend that139// changes LDS is PromoteAlloca and the changes it makes do not conflict.140//141// Addresses are written to constant global arrays based on the same metadata.142//143// The backend lowers LDS variables in the order of traversal of the function.144// This is at odds with the deterministic layout required. The workaround is to145// allocate the fixed-address variables immediately upon starting the function146// where they can be placed as intended. This requires a means of mapping from147// the function to the variables that it allocates. For the module scope lds,148// this is via metadata indicating whether the variable is not required. If a149// pass deletes that metadata, a fatal error on disagreement with the absolute150// symbol metadata will occur. For kernel scope and dynamic, this is by _name_151// correspondence between the function and the variable. It requires the152// kernel to have a name (which is only a limitation for tests in practice) and153// for nothing to rename the corresponding symbols. This is a hazard if the pass154// is run multiple times during debugging. Alternative schemes considered all155// involve bespoke metadata.156//157// If the name correspondence can be replaced, multiple distinct kernels that158// have the same memory layout can map to the same kernel id (as the address159// itself is handled by the absolute symbol metadata) and that will allow more160// uses of the "kernel" style faster lowering and reduce the size of the lookup161// tables.162//163// There is a test that checks this does not fire for a graphics shader. This164// lowering is expected to work for graphics if the isKernel test is changed.165//166// The current markUsedByKernel is sufficient for PromoteAlloca but is elided167// before codegen. Replacing this with an equivalent intrinsic which lasts until168// shortly after the machine function lowering of LDS would help break the name169// mapping. The other part needed is probably to amend PromoteAlloca to embed170// the LDS variables it creates in the same struct created here. That avoids the171// current hazard where a PromoteAlloca LDS variable might be allocated before172// the kernel scope (and thus error on the address check). Given a new invariant173// that no LDS variables exist outside of the structs managed here, and an174// intrinsic that lasts until after the LDS frame lowering, it should be175// possible to drop the name mapping and fold equivalent memory layouts.176//177//===----------------------------------------------------------------------===//178 179#include "AMDGPU.h"180#include "AMDGPUMemoryUtils.h"181#include "AMDGPUTargetMachine.h"182#include "Utils/AMDGPUBaseInfo.h"183#include "llvm/ADT/BitVector.h"184#include "llvm/ADT/DenseMap.h"185#include "llvm/ADT/DenseSet.h"186#include "llvm/ADT/STLExtras.h"187#include "llvm/ADT/SetOperations.h"188#include "llvm/Analysis/CallGraph.h"189#include "llvm/Analysis/ScopedNoAliasAA.h"190#include "llvm/CodeGen/TargetPassConfig.h"191#include "llvm/IR/Constants.h"192#include "llvm/IR/DerivedTypes.h"193#include "llvm/IR/Dominators.h"194#include "llvm/IR/IRBuilder.h"195#include "llvm/IR/InlineAsm.h"196#include "llvm/IR/Instructions.h"197#include "llvm/IR/IntrinsicsAMDGPU.h"198#include "llvm/IR/MDBuilder.h"199#include "llvm/IR/ReplaceConstant.h"200#include "llvm/InitializePasses.h"201#include "llvm/Pass.h"202#include "llvm/Support/CommandLine.h"203#include "llvm/Support/Debug.h"204#include "llvm/Support/Format.h"205#include "llvm/Support/OptimizedStructLayout.h"206#include "llvm/Support/raw_ostream.h"207#include "llvm/Transforms/Utils/BasicBlockUtils.h"208#include "llvm/Transforms/Utils/ModuleUtils.h"209 210#include <vector>211 212#include <cstdio>213 214#define DEBUG_TYPE "amdgpu-lower-module-lds"215 216using namespace llvm;217using namespace AMDGPU;218 219namespace {220 221cl::opt<bool> SuperAlignLDSGlobals(222 "amdgpu-super-align-lds-globals",223 cl::desc("Increase alignment of LDS if it is not on align boundary"),224 cl::init(true), cl::Hidden);225 226enum class LoweringKind { module, table, kernel, hybrid };227cl::opt<LoweringKind> LoweringKindLoc(228 "amdgpu-lower-module-lds-strategy",229 cl::desc("Specify lowering strategy for function LDS access:"), cl::Hidden,230 cl::init(LoweringKind::hybrid),231 cl::values(232 clEnumValN(LoweringKind::table, "table", "Lower via table lookup"),233 clEnumValN(LoweringKind::module, "module", "Lower via module struct"),234 clEnumValN(235 LoweringKind::kernel, "kernel",236 "Lower variables reachable from one kernel, otherwise abort"),237 clEnumValN(LoweringKind::hybrid, "hybrid",238 "Lower via mixture of above strategies")));239 240template <typename T> std::vector<T> sortByName(std::vector<T> &&V) {241 llvm::sort(V, [](const auto *L, const auto *R) {242 return L->getName() < R->getName();243 });244 return {std::move(V)};245}246 247class AMDGPULowerModuleLDS {248 const AMDGPUTargetMachine &TM;249 250 static void251 removeLocalVarsFromUsedLists(Module &M,252 const DenseSet<GlobalVariable *> &LocalVars) {253 // The verifier rejects used lists containing an inttoptr of a constant254 // so remove the variables from these lists before replaceAllUsesWith255 SmallPtrSet<Constant *, 8> LocalVarsSet;256 for (GlobalVariable *LocalVar : LocalVars)257 LocalVarsSet.insert(cast<Constant>(LocalVar->stripPointerCasts()));258 259 removeFromUsedLists(260 M, [&LocalVarsSet](Constant *C) { return LocalVarsSet.count(C); });261 262 for (GlobalVariable *LocalVar : LocalVars)263 LocalVar->removeDeadConstantUsers();264 }265 266 static void markUsedByKernel(Function *Func, GlobalVariable *SGV) {267 // The llvm.amdgcn.module.lds instance is implicitly used by all kernels268 // that might call a function which accesses a field within it. This is269 // presently approximated to 'all kernels' if there are any such functions270 // in the module. This implicit use is redefined as an explicit use here so271 // that later passes, specifically PromoteAlloca, account for the required272 // memory without any knowledge of this transform.273 274 // An operand bundle on llvm.donothing works because the call instruction275 // survives until after the last pass that needs to account for LDS. It is276 // better than inline asm as the latter survives until the end of codegen. A277 // totally robust solution would be a function with the same semantics as278 // llvm.donothing that takes a pointer to the instance and is lowered to a279 // no-op after LDS is allocated, but that is not presently necessary.280 281 // This intrinsic is eliminated shortly before instruction selection. It282 // does not suffice to indicate to ISel that a given global which is not283 // immediately used by the kernel must still be allocated by it. An284 // equivalent target specific intrinsic which lasts until immediately after285 // codegen would suffice for that, but one would still need to ensure that286 // the variables are allocated in the anticipated order.287 BasicBlock *Entry = &Func->getEntryBlock();288 IRBuilder<> Builder(Entry, Entry->getFirstNonPHIIt());289 290 Function *Decl = Intrinsic::getOrInsertDeclaration(291 Func->getParent(), Intrinsic::donothing, {});292 293 Value *UseInstance[1] = {294 Builder.CreateConstInBoundsGEP1_32(SGV->getValueType(), SGV, 0)};295 296 Builder.CreateCall(297 Decl, {}, {OperandBundleDefT<Value *>("ExplicitUse", UseInstance)});298 }299 300public:301 AMDGPULowerModuleLDS(const AMDGPUTargetMachine &TM_) : TM(TM_) {}302 303 struct LDSVariableReplacement {304 GlobalVariable *SGV = nullptr;305 DenseMap<GlobalVariable *, Constant *> LDSVarsToConstantGEP;306 };307 308 // remap from lds global to a constantexpr gep to where it has been moved to309 // for each kernel310 // an array with an element for each kernel containing where the corresponding311 // variable was remapped to312 313 static Constant *getAddressesOfVariablesInKernel(314 LLVMContext &Ctx, ArrayRef<GlobalVariable *> Variables,315 const DenseMap<GlobalVariable *, Constant *> &LDSVarsToConstantGEP) {316 // Create a ConstantArray containing the address of each Variable within the317 // kernel corresponding to LDSVarsToConstantGEP, or poison if that kernel318 // does not allocate it319 // TODO: Drop the ptrtoint conversion320 321 Type *I32 = Type::getInt32Ty(Ctx);322 323 ArrayType *KernelOffsetsType = ArrayType::get(I32, Variables.size());324 325 SmallVector<Constant *> Elements;326 for (GlobalVariable *GV : Variables) {327 auto ConstantGepIt = LDSVarsToConstantGEP.find(GV);328 if (ConstantGepIt != LDSVarsToConstantGEP.end()) {329 auto *elt = ConstantExpr::getPtrToInt(ConstantGepIt->second, I32);330 Elements.push_back(elt);331 } else {332 Elements.push_back(PoisonValue::get(I32));333 }334 }335 return ConstantArray::get(KernelOffsetsType, Elements);336 }337 338 static GlobalVariable *buildLookupTable(339 Module &M, ArrayRef<GlobalVariable *> Variables,340 ArrayRef<Function *> kernels,341 DenseMap<Function *, LDSVariableReplacement> &KernelToReplacement) {342 if (Variables.empty()) {343 return nullptr;344 }345 LLVMContext &Ctx = M.getContext();346 347 const size_t NumberVariables = Variables.size();348 const size_t NumberKernels = kernels.size();349 350 ArrayType *KernelOffsetsType =351 ArrayType::get(Type::getInt32Ty(Ctx), NumberVariables);352 353 ArrayType *AllKernelsOffsetsType =354 ArrayType::get(KernelOffsetsType, NumberKernels);355 356 Constant *Missing = PoisonValue::get(KernelOffsetsType);357 std::vector<Constant *> overallConstantExprElts(NumberKernels);358 for (size_t i = 0; i < NumberKernels; i++) {359 auto Replacement = KernelToReplacement.find(kernels[i]);360 overallConstantExprElts[i] =361 (Replacement == KernelToReplacement.end())362 ? Missing363 : getAddressesOfVariablesInKernel(364 Ctx, Variables, Replacement->second.LDSVarsToConstantGEP);365 }366 367 Constant *init =368 ConstantArray::get(AllKernelsOffsetsType, overallConstantExprElts);369 370 return new GlobalVariable(371 M, AllKernelsOffsetsType, true, GlobalValue::InternalLinkage, init,372 "llvm.amdgcn.lds.offset.table", nullptr, GlobalValue::NotThreadLocal,373 AMDGPUAS::CONSTANT_ADDRESS);374 }375 376 void replaceUseWithTableLookup(Module &M, IRBuilder<> &Builder,377 GlobalVariable *LookupTable,378 GlobalVariable *GV, Use &U,379 Value *OptionalIndex) {380 // Table is a constant array of the same length as OrderedKernels381 LLVMContext &Ctx = M.getContext();382 Type *I32 = Type::getInt32Ty(Ctx);383 auto *I = cast<Instruction>(U.getUser());384 385 Value *tableKernelIndex = getTableLookupKernelIndex(M, I->getFunction());386 387 if (auto *Phi = dyn_cast<PHINode>(I)) {388 BasicBlock *BB = Phi->getIncomingBlock(U);389 Builder.SetInsertPoint(&(*(BB->getFirstInsertionPt())));390 } else {391 Builder.SetInsertPoint(I);392 }393 394 SmallVector<Value *, 3> GEPIdx = {395 ConstantInt::get(I32, 0),396 tableKernelIndex,397 };398 if (OptionalIndex)399 GEPIdx.push_back(OptionalIndex);400 401 Value *Address = Builder.CreateInBoundsGEP(402 LookupTable->getValueType(), LookupTable, GEPIdx, GV->getName());403 404 Value *loaded = Builder.CreateLoad(I32, Address);405 406 Value *replacement =407 Builder.CreateIntToPtr(loaded, GV->getType(), GV->getName());408 409 U.set(replacement);410 }411 412 void replaceUsesInInstructionsWithTableLookup(413 Module &M, ArrayRef<GlobalVariable *> ModuleScopeVariables,414 GlobalVariable *LookupTable) {415 416 LLVMContext &Ctx = M.getContext();417 IRBuilder<> Builder(Ctx);418 Type *I32 = Type::getInt32Ty(Ctx);419 420 for (size_t Index = 0; Index < ModuleScopeVariables.size(); Index++) {421 auto *GV = ModuleScopeVariables[Index];422 423 for (Use &U : make_early_inc_range(GV->uses())) {424 auto *I = dyn_cast<Instruction>(U.getUser());425 if (!I)426 continue;427 428 replaceUseWithTableLookup(M, Builder, LookupTable, GV, U,429 ConstantInt::get(I32, Index));430 }431 }432 }433 434 static DenseSet<Function *> kernelsThatIndirectlyAccessAnyOfPassedVariables(435 Module &M, LDSUsesInfoTy &LDSUsesInfo,436 DenseSet<GlobalVariable *> const &VariableSet) {437 438 DenseSet<Function *> KernelSet;439 440 if (VariableSet.empty())441 return KernelSet;442 443 for (Function &Func : M.functions()) {444 if (Func.isDeclaration() || !isKernel(Func))445 continue;446 for (GlobalVariable *GV : LDSUsesInfo.indirect_access[&Func]) {447 if (VariableSet.contains(GV)) {448 KernelSet.insert(&Func);449 break;450 }451 }452 }453 454 return KernelSet;455 }456 457 static GlobalVariable *458 chooseBestVariableForModuleStrategy(const DataLayout &DL,459 VariableFunctionMap &LDSVars) {460 // Find the global variable with the most indirect uses from kernels461 462 struct CandidateTy {463 GlobalVariable *GV = nullptr;464 size_t UserCount = 0;465 size_t Size = 0;466 467 CandidateTy() = default;468 469 CandidateTy(GlobalVariable *GV, uint64_t UserCount, uint64_t AllocSize)470 : GV(GV), UserCount(UserCount), Size(AllocSize) {}471 472 bool operator<(const CandidateTy &Other) const {473 // Fewer users makes module scope variable less attractive474 if (UserCount < Other.UserCount) {475 return true;476 }477 if (UserCount > Other.UserCount) {478 return false;479 }480 481 // Bigger makes module scope variable less attractive482 if (Size < Other.Size) {483 return false;484 }485 486 if (Size > Other.Size) {487 return true;488 }489 490 // Arbitrary but consistent491 return GV->getName() < Other.GV->getName();492 }493 };494 495 CandidateTy MostUsed;496 497 for (auto &K : LDSVars) {498 GlobalVariable *GV = K.first;499 if (K.second.size() <= 1) {500 // A variable reachable by only one kernel is best lowered with kernel501 // strategy502 continue;503 }504 CandidateTy Candidate(505 GV, K.second.size(),506 DL.getTypeAllocSize(GV->getValueType()).getFixedValue());507 if (MostUsed < Candidate)508 MostUsed = Candidate;509 }510 511 return MostUsed.GV;512 }513 514 static void recordLDSAbsoluteAddress(Module *M, GlobalVariable *GV,515 uint32_t Address) {516 // Write the specified address into metadata where it can be retrieved by517 // the assembler. Format is a half open range, [Address Address+1)518 LLVMContext &Ctx = M->getContext();519 auto *IntTy =520 M->getDataLayout().getIntPtrType(Ctx, AMDGPUAS::LOCAL_ADDRESS);521 auto *MinC = ConstantAsMetadata::get(ConstantInt::get(IntTy, Address));522 auto *MaxC = ConstantAsMetadata::get(ConstantInt::get(IntTy, Address + 1));523 GV->setMetadata(LLVMContext::MD_absolute_symbol,524 MDNode::get(Ctx, {MinC, MaxC}));525 }526 527 DenseMap<Function *, Value *> tableKernelIndexCache;528 Value *getTableLookupKernelIndex(Module &M, Function *F) {529 // Accesses from a function use the amdgcn_lds_kernel_id intrinsic which530 // lowers to a read from a live in register. Emit it once in the entry531 // block to spare deduplicating it later.532 auto [It, Inserted] = tableKernelIndexCache.try_emplace(F);533 if (Inserted) {534 auto InsertAt = F->getEntryBlock().getFirstNonPHIOrDbgOrAlloca();535 IRBuilder<> Builder(&*InsertAt);536 537 It->second = Builder.CreateIntrinsic(Intrinsic::amdgcn_lds_kernel_id, {});538 }539 540 return It->second;541 }542 543 static std::vector<Function *> assignLDSKernelIDToEachKernel(544 Module *M, DenseSet<Function *> const &KernelsThatAllocateTableLDS,545 DenseSet<Function *> const &KernelsThatIndirectlyAllocateDynamicLDS) {546 // Associate kernels in the set with an arbitrary but reproducible order and547 // annotate them with that order in metadata. This metadata is recognised by548 // the backend and lowered to a SGPR which can be read from using549 // amdgcn_lds_kernel_id.550 551 std::vector<Function *> OrderedKernels;552 if (!KernelsThatAllocateTableLDS.empty() ||553 !KernelsThatIndirectlyAllocateDynamicLDS.empty()) {554 555 for (Function &Func : M->functions()) {556 if (Func.isDeclaration())557 continue;558 if (!isKernel(Func))559 continue;560 561 if (KernelsThatAllocateTableLDS.contains(&Func) ||562 KernelsThatIndirectlyAllocateDynamicLDS.contains(&Func)) {563 assert(Func.hasName()); // else fatal error earlier564 OrderedKernels.push_back(&Func);565 }566 }567 568 // Put them in an arbitrary but reproducible order569 OrderedKernels = sortByName(std::move(OrderedKernels));570 571 // Annotate the kernels with their order in this vector572 LLVMContext &Ctx = M->getContext();573 IRBuilder<> Builder(Ctx);574 575 if (OrderedKernels.size() > UINT32_MAX) {576 // 32 bit keeps it in one SGPR. > 2**32 kernels won't fit on the GPU577 reportFatalUsageError("unimplemented LDS lowering for > 2**32 kernels");578 }579 580 for (size_t i = 0; i < OrderedKernels.size(); i++) {581 Metadata *AttrMDArgs[1] = {582 ConstantAsMetadata::get(Builder.getInt32(i)),583 };584 OrderedKernels[i]->setMetadata("llvm.amdgcn.lds.kernel.id",585 MDNode::get(Ctx, AttrMDArgs));586 }587 }588 return OrderedKernels;589 }590 591 static void partitionVariablesIntoIndirectStrategies(592 Module &M, LDSUsesInfoTy const &LDSUsesInfo,593 VariableFunctionMap &LDSToKernelsThatNeedToAccessItIndirectly,594 DenseSet<GlobalVariable *> &ModuleScopeVariables,595 DenseSet<GlobalVariable *> &TableLookupVariables,596 DenseSet<GlobalVariable *> &KernelAccessVariables,597 DenseSet<GlobalVariable *> &DynamicVariables) {598 599 GlobalVariable *HybridModuleRoot =600 LoweringKindLoc != LoweringKind::hybrid601 ? nullptr602 : chooseBestVariableForModuleStrategy(603 M.getDataLayout(), LDSToKernelsThatNeedToAccessItIndirectly);604 605 DenseSet<Function *> const EmptySet;606 DenseSet<Function *> const &HybridModuleRootKernels =607 HybridModuleRoot608 ? LDSToKernelsThatNeedToAccessItIndirectly[HybridModuleRoot]609 : EmptySet;610 611 for (auto &K : LDSToKernelsThatNeedToAccessItIndirectly) {612 // Each iteration of this loop assigns exactly one global variable to613 // exactly one of the implementation strategies.614 615 GlobalVariable *GV = K.first;616 assert(AMDGPU::isLDSVariableToLower(*GV));617 assert(K.second.size() != 0);618 619 if (AMDGPU::isDynamicLDS(*GV)) {620 DynamicVariables.insert(GV);621 continue;622 }623 624 switch (LoweringKindLoc) {625 case LoweringKind::module:626 ModuleScopeVariables.insert(GV);627 break;628 629 case LoweringKind::table:630 TableLookupVariables.insert(GV);631 break;632 633 case LoweringKind::kernel:634 if (K.second.size() == 1) {635 KernelAccessVariables.insert(GV);636 } else {637 // FIXME: This should use DiagnosticInfo638 reportFatalUsageError(639 "cannot lower LDS '" + GV->getName() +640 "' to kernel access as it is reachable from multiple kernels");641 }642 break;643 644 case LoweringKind::hybrid: {645 if (GV == HybridModuleRoot) {646 assert(K.second.size() != 1);647 ModuleScopeVariables.insert(GV);648 } else if (K.second.size() == 1) {649 KernelAccessVariables.insert(GV);650 } else if (K.second == HybridModuleRootKernels) {651 ModuleScopeVariables.insert(GV);652 } else {653 TableLookupVariables.insert(GV);654 }655 break;656 }657 }658 }659 660 // All LDS variables accessed indirectly have now been partitioned into661 // the distinct lowering strategies.662 assert(ModuleScopeVariables.size() + TableLookupVariables.size() +663 KernelAccessVariables.size() + DynamicVariables.size() ==664 LDSToKernelsThatNeedToAccessItIndirectly.size());665 }666 667 static GlobalVariable *lowerModuleScopeStructVariables(668 Module &M, DenseSet<GlobalVariable *> const &ModuleScopeVariables,669 DenseSet<Function *> const &KernelsThatAllocateModuleLDS) {670 // Create a struct to hold the ModuleScopeVariables671 // Replace all uses of those variables from non-kernel functions with the672 // new struct instance Replace only the uses from kernel functions that will673 // allocate this instance. That is a space optimisation - kernels that use a674 // subset of the module scope struct and do not need to allocate it for675 // indirect calls will only allocate the subset they use (they do so as part676 // of the per-kernel lowering).677 if (ModuleScopeVariables.empty()) {678 return nullptr;679 }680 681 LLVMContext &Ctx = M.getContext();682 683 LDSVariableReplacement ModuleScopeReplacement =684 createLDSVariableReplacement(M, "llvm.amdgcn.module.lds",685 ModuleScopeVariables);686 687 appendToCompilerUsed(M, {static_cast<GlobalValue *>(688 ConstantExpr::getPointerBitCastOrAddrSpaceCast(689 cast<Constant>(ModuleScopeReplacement.SGV),690 PointerType::getUnqual(Ctx)))});691 692 // module.lds will be allocated at zero in any kernel that allocates it693 recordLDSAbsoluteAddress(&M, ModuleScopeReplacement.SGV, 0);694 695 // historic696 removeLocalVarsFromUsedLists(M, ModuleScopeVariables);697 698 // Replace all uses of module scope variable from non-kernel functions699 replaceLDSVariablesWithStruct(700 M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) {701 Instruction *I = dyn_cast<Instruction>(U.getUser());702 if (!I) {703 return false;704 }705 Function *F = I->getFunction();706 return !isKernel(*F);707 });708 709 // Replace uses of module scope variable from kernel functions that710 // allocate the module scope variable, otherwise leave them unchanged711 // Record on each kernel whether the module scope global is used by it712 713 for (Function &Func : M.functions()) {714 if (Func.isDeclaration() || !isKernel(Func))715 continue;716 717 if (KernelsThatAllocateModuleLDS.contains(&Func)) {718 replaceLDSVariablesWithStruct(719 M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) {720 Instruction *I = dyn_cast<Instruction>(U.getUser());721 if (!I) {722 return false;723 }724 Function *F = I->getFunction();725 return F == &Func;726 });727 728 markUsedByKernel(&Func, ModuleScopeReplacement.SGV);729 }730 }731 732 return ModuleScopeReplacement.SGV;733 }734 735 static DenseMap<Function *, LDSVariableReplacement>736 lowerKernelScopeStructVariables(737 Module &M, LDSUsesInfoTy &LDSUsesInfo,738 DenseSet<GlobalVariable *> const &ModuleScopeVariables,739 DenseSet<Function *> const &KernelsThatAllocateModuleLDS,740 GlobalVariable *MaybeModuleScopeStruct) {741 742 // Create a struct for each kernel for the non-module-scope variables.743 744 DenseMap<Function *, LDSVariableReplacement> KernelToReplacement;745 for (Function &Func : M.functions()) {746 if (Func.isDeclaration() || !isKernel(Func))747 continue;748 749 DenseSet<GlobalVariable *> KernelUsedVariables;750 // Allocating variables that are used directly in this struct to get751 // alignment aware allocation and predictable frame size.752 for (auto &v : LDSUsesInfo.direct_access[&Func]) {753 if (!AMDGPU::isDynamicLDS(*v)) {754 KernelUsedVariables.insert(v);755 }756 }757 758 // Allocating variables that are accessed indirectly so that a lookup of759 // this struct instance can find them from nested functions.760 for (auto &v : LDSUsesInfo.indirect_access[&Func]) {761 if (!AMDGPU::isDynamicLDS(*v)) {762 KernelUsedVariables.insert(v);763 }764 }765 766 // Variables allocated in module lds must all resolve to that struct,767 // not to the per-kernel instance.768 if (KernelsThatAllocateModuleLDS.contains(&Func)) {769 for (GlobalVariable *v : ModuleScopeVariables) {770 KernelUsedVariables.erase(v);771 }772 }773 774 if (KernelUsedVariables.empty()) {775 // Either used no LDS, or the LDS it used was all in the module struct776 // or dynamically sized777 continue;778 }779 780 // The association between kernel function and LDS struct is done by781 // symbol name, which only works if the function in question has a782 // name This is not expected to be a problem in practice as kernels783 // are called by name making anonymous ones (which are named by the784 // backend) difficult to use. This does mean that llvm test cases need785 // to name the kernels.786 if (!Func.hasName()) {787 reportFatalUsageError("anonymous kernels cannot use LDS variables");788 }789 790 std::string VarName =791 (Twine("llvm.amdgcn.kernel.") + Func.getName() + ".lds").str();792 793 auto Replacement =794 createLDSVariableReplacement(M, VarName, KernelUsedVariables);795 796 // If any indirect uses, create a direct use to ensure allocation797 // TODO: Simpler to unconditionally mark used but that regresses798 // codegen in test/CodeGen/AMDGPU/noclobber-barrier.ll799 auto Accesses = LDSUsesInfo.indirect_access.find(&Func);800 if ((Accesses != LDSUsesInfo.indirect_access.end()) &&801 !Accesses->second.empty())802 markUsedByKernel(&Func, Replacement.SGV);803 804 // remove preserves existing codegen805 removeLocalVarsFromUsedLists(M, KernelUsedVariables);806 KernelToReplacement[&Func] = Replacement;807 808 // Rewrite uses within kernel to the new struct809 replaceLDSVariablesWithStruct(810 M, KernelUsedVariables, Replacement, [&Func](Use &U) {811 Instruction *I = dyn_cast<Instruction>(U.getUser());812 return I && I->getFunction() == &Func;813 });814 }815 return KernelToReplacement;816 }817 818 static GlobalVariable *819 buildRepresentativeDynamicLDSInstance(Module &M, LDSUsesInfoTy &LDSUsesInfo,820 Function *func) {821 // Create a dynamic lds variable with a name associated with the passed822 // function that has the maximum alignment of any dynamic lds variable823 // reachable from this kernel. Dynamic LDS is allocated after the static LDS824 // allocation, possibly after alignment padding. The representative variable825 // created here has the maximum alignment of any other dynamic variable826 // reachable by that kernel. All dynamic LDS variables are allocated at the827 // same address in each kernel in order to provide the documented aliasing828 // semantics. Setting the alignment here allows this IR pass to accurately829 // predict the exact constant at which it will be allocated.830 831 assert(isKernel(*func));832 833 LLVMContext &Ctx = M.getContext();834 const DataLayout &DL = M.getDataLayout();835 Align MaxDynamicAlignment(1);836 837 auto UpdateMaxAlignment = [&MaxDynamicAlignment, &DL](GlobalVariable *GV) {838 if (AMDGPU::isDynamicLDS(*GV)) {839 MaxDynamicAlignment =840 std::max(MaxDynamicAlignment, AMDGPU::getAlign(DL, GV));841 }842 };843 844 for (GlobalVariable *GV : LDSUsesInfo.indirect_access[func]) {845 UpdateMaxAlignment(GV);846 }847 848 for (GlobalVariable *GV : LDSUsesInfo.direct_access[func]) {849 UpdateMaxAlignment(GV);850 }851 852 assert(func->hasName()); // Checked by caller853 auto *emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);854 GlobalVariable *N = new GlobalVariable(855 M, emptyCharArray, false, GlobalValue::ExternalLinkage, nullptr,856 Twine("llvm.amdgcn." + func->getName() + ".dynlds"), nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS,857 false);858 N->setAlignment(MaxDynamicAlignment);859 860 assert(AMDGPU::isDynamicLDS(*N));861 return N;862 }863 864 DenseMap<Function *, GlobalVariable *> lowerDynamicLDSVariables(865 Module &M, LDSUsesInfoTy &LDSUsesInfo,866 DenseSet<Function *> const &KernelsThatIndirectlyAllocateDynamicLDS,867 DenseSet<GlobalVariable *> const &DynamicVariables,868 std::vector<Function *> const &OrderedKernels) {869 DenseMap<Function *, GlobalVariable *> KernelToCreatedDynamicLDS;870 if (!KernelsThatIndirectlyAllocateDynamicLDS.empty()) {871 LLVMContext &Ctx = M.getContext();872 IRBuilder<> Builder(Ctx);873 Type *I32 = Type::getInt32Ty(Ctx);874 875 std::vector<Constant *> newDynamicLDS;876 877 // Table is built in the same order as OrderedKernels878 for (auto &func : OrderedKernels) {879 880 if (KernelsThatIndirectlyAllocateDynamicLDS.contains(func)) {881 assert(isKernel(*func));882 if (!func->hasName()) {883 reportFatalUsageError("anonymous kernels cannot use LDS variables");884 }885 886 GlobalVariable *N =887 buildRepresentativeDynamicLDSInstance(M, LDSUsesInfo, func);888 889 KernelToCreatedDynamicLDS[func] = N;890 891 markUsedByKernel(func, N);892 893 auto *emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);894 auto *GEP = ConstantExpr::getGetElementPtr(895 emptyCharArray, N, ConstantInt::get(I32, 0), true);896 newDynamicLDS.push_back(ConstantExpr::getPtrToInt(GEP, I32));897 } else {898 newDynamicLDS.push_back(PoisonValue::get(I32));899 }900 }901 assert(OrderedKernels.size() == newDynamicLDS.size());902 903 ArrayType *t = ArrayType::get(I32, newDynamicLDS.size());904 Constant *init = ConstantArray::get(t, newDynamicLDS);905 GlobalVariable *table = new GlobalVariable(906 M, t, true, GlobalValue::InternalLinkage, init,907 "llvm.amdgcn.dynlds.offset.table", nullptr,908 GlobalValue::NotThreadLocal, AMDGPUAS::CONSTANT_ADDRESS);909 910 for (GlobalVariable *GV : DynamicVariables) {911 for (Use &U : make_early_inc_range(GV->uses())) {912 auto *I = dyn_cast<Instruction>(U.getUser());913 if (!I)914 continue;915 if (isKernel(*I->getFunction()))916 continue;917 918 replaceUseWithTableLookup(M, Builder, table, GV, U, nullptr);919 }920 }921 }922 return KernelToCreatedDynamicLDS;923 }924 925 bool runOnModule(Module &M) {926 CallGraph CG = CallGraph(M);927 bool Changed = superAlignLDSGlobals(M);928 929 Changed |= eliminateConstantExprUsesOfLDSFromAllInstructions(M);930 931 Changed = true; // todo: narrow this down932 933 // For each kernel, what variables does it access directly or through934 // callees935 LDSUsesInfoTy LDSUsesInfo = getTransitiveUsesOfLDS(CG, M);936 937 // For each variable accessed through callees, which kernels access it938 VariableFunctionMap LDSToKernelsThatNeedToAccessItIndirectly;939 for (auto &K : LDSUsesInfo.indirect_access) {940 Function *F = K.first;941 assert(isKernel(*F));942 for (GlobalVariable *GV : K.second) {943 LDSToKernelsThatNeedToAccessItIndirectly[GV].insert(F);944 }945 }946 947 // Partition variables accessed indirectly into the different strategies948 DenseSet<GlobalVariable *> ModuleScopeVariables;949 DenseSet<GlobalVariable *> TableLookupVariables;950 DenseSet<GlobalVariable *> KernelAccessVariables;951 DenseSet<GlobalVariable *> DynamicVariables;952 partitionVariablesIntoIndirectStrategies(953 M, LDSUsesInfo, LDSToKernelsThatNeedToAccessItIndirectly,954 ModuleScopeVariables, TableLookupVariables, KernelAccessVariables,955 DynamicVariables);956 957 // If the kernel accesses a variable that is going to be stored in the958 // module instance through a call then that kernel needs to allocate the959 // module instance960 const DenseSet<Function *> KernelsThatAllocateModuleLDS =961 kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo,962 ModuleScopeVariables);963 const DenseSet<Function *> KernelsThatAllocateTableLDS =964 kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo,965 TableLookupVariables);966 967 const DenseSet<Function *> KernelsThatIndirectlyAllocateDynamicLDS =968 kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo,969 DynamicVariables);970 971 GlobalVariable *MaybeModuleScopeStruct = lowerModuleScopeStructVariables(972 M, ModuleScopeVariables, KernelsThatAllocateModuleLDS);973 974 DenseMap<Function *, LDSVariableReplacement> KernelToReplacement =975 lowerKernelScopeStructVariables(M, LDSUsesInfo, ModuleScopeVariables,976 KernelsThatAllocateModuleLDS,977 MaybeModuleScopeStruct);978 979 // Lower zero cost accesses to the kernel instances just created980 for (auto &GV : KernelAccessVariables) {981 auto &funcs = LDSToKernelsThatNeedToAccessItIndirectly[GV];982 assert(funcs.size() == 1); // Only one kernel can access it983 LDSVariableReplacement Replacement =984 KernelToReplacement[*(funcs.begin())];985 986 DenseSet<GlobalVariable *> Vec;987 Vec.insert(GV);988 989 replaceLDSVariablesWithStruct(M, Vec, Replacement, [](Use &U) {990 return isa<Instruction>(U.getUser());991 });992 }993 994 // The ith element of this vector is kernel id i995 std::vector<Function *> OrderedKernels =996 assignLDSKernelIDToEachKernel(&M, KernelsThatAllocateTableLDS,997 KernelsThatIndirectlyAllocateDynamicLDS);998 999 if (!KernelsThatAllocateTableLDS.empty()) {1000 LLVMContext &Ctx = M.getContext();1001 IRBuilder<> Builder(Ctx);1002 1003 // The order must be consistent between lookup table and accesses to1004 // lookup table1005 auto TableLookupVariablesOrdered =1006 sortByName(std::vector<GlobalVariable *>(TableLookupVariables.begin(),1007 TableLookupVariables.end()));1008 1009 GlobalVariable *LookupTable = buildLookupTable(1010 M, TableLookupVariablesOrdered, OrderedKernels, KernelToReplacement);1011 replaceUsesInInstructionsWithTableLookup(M, TableLookupVariablesOrdered,1012 LookupTable);1013 }1014 1015 DenseMap<Function *, GlobalVariable *> KernelToCreatedDynamicLDS =1016 lowerDynamicLDSVariables(M, LDSUsesInfo,1017 KernelsThatIndirectlyAllocateDynamicLDS,1018 DynamicVariables, OrderedKernels);1019 1020 // Strip amdgpu-no-lds-kernel-id from all functions reachable from the1021 // kernel. We may have inferred this wasn't used prior to the pass.1022 // TODO: We could filter out subgraphs that do not access LDS globals.1023 for (auto *KernelSet : {&KernelsThatIndirectlyAllocateDynamicLDS,1024 &KernelsThatAllocateTableLDS})1025 for (Function *F : *KernelSet)1026 removeFnAttrFromReachable(CG, F, {"amdgpu-no-lds-kernel-id"});1027 1028 // All kernel frames have been allocated. Calculate and record the1029 // addresses.1030 {1031 const DataLayout &DL = M.getDataLayout();1032 1033 for (Function &Func : M.functions()) {1034 if (Func.isDeclaration() || !isKernel(Func))1035 continue;1036 1037 // All three of these are optional. The first variable is allocated at1038 // zero. They are allocated by AMDGPUMachineFunction as one block.1039 // Layout:1040 //{1041 // module.lds1042 // alignment padding1043 // kernel instance1044 // alignment padding1045 // dynamic lds variables1046 //}1047 1048 const bool AllocateModuleScopeStruct =1049 MaybeModuleScopeStruct &&1050 KernelsThatAllocateModuleLDS.contains(&Func);1051 1052 auto Replacement = KernelToReplacement.find(&Func);1053 const bool AllocateKernelScopeStruct =1054 Replacement != KernelToReplacement.end();1055 1056 const bool AllocateDynamicVariable =1057 KernelToCreatedDynamicLDS.contains(&Func);1058 1059 uint32_t Offset = 0;1060 1061 if (AllocateModuleScopeStruct) {1062 // Allocated at zero, recorded once on construction, not once per1063 // kernel1064 Offset += DL.getTypeAllocSize(MaybeModuleScopeStruct->getValueType());1065 }1066 1067 if (AllocateKernelScopeStruct) {1068 GlobalVariable *KernelStruct = Replacement->second.SGV;1069 Offset = alignTo(Offset, AMDGPU::getAlign(DL, KernelStruct));1070 recordLDSAbsoluteAddress(&M, KernelStruct, Offset);1071 Offset += DL.getTypeAllocSize(KernelStruct->getValueType());1072 }1073 1074 // If there is dynamic allocation, the alignment needed is included in1075 // the static frame size. There may be no reference to the dynamic1076 // variable in the kernel itself, so without including it here, that1077 // alignment padding could be missed.1078 if (AllocateDynamicVariable) {1079 GlobalVariable *DynamicVariable = KernelToCreatedDynamicLDS[&Func];1080 Offset = alignTo(Offset, AMDGPU::getAlign(DL, DynamicVariable));1081 recordLDSAbsoluteAddress(&M, DynamicVariable, Offset);1082 }1083 1084 if (Offset != 0) {1085 (void)TM; // TODO: Account for target maximum LDS1086 std::string Buffer;1087 raw_string_ostream SS{Buffer};1088 SS << format("%u", Offset);1089 1090 // Instead of explicitly marking kernels that access dynamic variables1091 // using special case metadata, annotate with min-lds == max-lds, i.e.1092 // that there is no more space available for allocating more static1093 // LDS variables. That is the right condition to prevent allocating1094 // more variables which would collide with the addresses assigned to1095 // dynamic variables.1096 if (AllocateDynamicVariable)1097 SS << format(",%u", Offset);1098 1099 Func.addFnAttr("amdgpu-lds-size", Buffer);1100 }1101 }1102 }1103 1104 for (auto &GV : make_early_inc_range(M.globals()))1105 if (AMDGPU::isLDSVariableToLower(GV)) {1106 // probably want to remove from used lists1107 GV.removeDeadConstantUsers();1108 if (GV.use_empty())1109 GV.eraseFromParent();1110 }1111 1112 return Changed;1113 }1114 1115private:1116 // Increase the alignment of LDS globals if necessary to maximise the chance1117 // that we can use aligned LDS instructions to access them.1118 static bool superAlignLDSGlobals(Module &M) {1119 const DataLayout &DL = M.getDataLayout();1120 bool Changed = false;1121 if (!SuperAlignLDSGlobals) {1122 return Changed;1123 }1124 1125 for (auto &GV : M.globals()) {1126 if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {1127 // Only changing alignment of LDS variables1128 continue;1129 }1130 if (!GV.hasInitializer()) {1131 // cuda/hip extern __shared__ variable, leave alignment alone1132 continue;1133 }1134 1135 if (GV.isAbsoluteSymbolRef()) {1136 // If the variable is already allocated, don't change the alignment1137 continue;1138 }1139 1140 Align Alignment = AMDGPU::getAlign(DL, &GV);1141 TypeSize GVSize = DL.getTypeAllocSize(GV.getValueType());1142 1143 if (GVSize > 8) {1144 // We might want to use a b96 or b128 load/store1145 Alignment = std::max(Alignment, Align(16));1146 } else if (GVSize > 4) {1147 // We might want to use a b64 load/store1148 Alignment = std::max(Alignment, Align(8));1149 } else if (GVSize > 2) {1150 // We might want to use a b32 load/store1151 Alignment = std::max(Alignment, Align(4));1152 } else if (GVSize > 1) {1153 // We might want to use a b16 load/store1154 Alignment = std::max(Alignment, Align(2));1155 }1156 1157 if (Alignment != AMDGPU::getAlign(DL, &GV)) {1158 Changed = true;1159 GV.setAlignment(Alignment);1160 }1161 }1162 return Changed;1163 }1164 1165 static LDSVariableReplacement createLDSVariableReplacement(1166 Module &M, std::string VarName,1167 DenseSet<GlobalVariable *> const &LDSVarsToTransform) {1168 // Create a struct instance containing LDSVarsToTransform and map from those1169 // variables to ConstantExprGEP1170 // Variables may be introduced to meet alignment requirements. No aliasing1171 // metadata is useful for these as they have no uses. Erased before return.1172 1173 LLVMContext &Ctx = M.getContext();1174 const DataLayout &DL = M.getDataLayout();1175 assert(!LDSVarsToTransform.empty());1176 1177 SmallVector<OptimizedStructLayoutField, 8> LayoutFields;1178 LayoutFields.reserve(LDSVarsToTransform.size());1179 {1180 // The order of fields in this struct depends on the order of1181 // variables in the argument which varies when changing how they1182 // are identified, leading to spurious test breakage.1183 auto Sorted = sortByName(std::vector<GlobalVariable *>(1184 LDSVarsToTransform.begin(), LDSVarsToTransform.end()));1185 1186 for (GlobalVariable *GV : Sorted) {1187 OptimizedStructLayoutField F(GV,1188 DL.getTypeAllocSize(GV->getValueType()),1189 AMDGPU::getAlign(DL, GV));1190 LayoutFields.emplace_back(F);1191 }1192 }1193 1194 performOptimizedStructLayout(LayoutFields);1195 1196 std::vector<GlobalVariable *> LocalVars;1197 BitVector IsPaddingField;1198 LocalVars.reserve(LDSVarsToTransform.size()); // will be at least this large1199 IsPaddingField.reserve(LDSVarsToTransform.size());1200 {1201 uint64_t CurrentOffset = 0;1202 for (auto &F : LayoutFields) {1203 GlobalVariable *FGV =1204 static_cast<GlobalVariable *>(const_cast<void *>(F.Id));1205 Align DataAlign = F.Alignment;1206 1207 uint64_t DataAlignV = DataAlign.value();1208 if (uint64_t Rem = CurrentOffset % DataAlignV) {1209 uint64_t Padding = DataAlignV - Rem;1210 1211 // Append an array of padding bytes to meet alignment requested1212 // Note (o + (a - (o % a)) ) % a == 01213 // (offset + Padding ) % align == 01214 1215 Type *ATy = ArrayType::get(Type::getInt8Ty(Ctx), Padding);1216 LocalVars.push_back(new GlobalVariable(1217 M, ATy, false, GlobalValue::InternalLinkage,1218 PoisonValue::get(ATy), "", nullptr, GlobalValue::NotThreadLocal,1219 AMDGPUAS::LOCAL_ADDRESS, false));1220 IsPaddingField.push_back(true);1221 CurrentOffset += Padding;1222 }1223 1224 LocalVars.push_back(FGV);1225 IsPaddingField.push_back(false);1226 CurrentOffset += F.Size;1227 }1228 }1229 1230 std::vector<Type *> LocalVarTypes;1231 LocalVarTypes.reserve(LocalVars.size());1232 std::transform(1233 LocalVars.cbegin(), LocalVars.cend(), std::back_inserter(LocalVarTypes),1234 [](const GlobalVariable *V) -> Type * { return V->getValueType(); });1235 1236 StructType *LDSTy = StructType::create(Ctx, LocalVarTypes, VarName + ".t");1237 1238 Align StructAlign = AMDGPU::getAlign(DL, LocalVars[0]);1239 1240 GlobalVariable *SGV = new GlobalVariable(1241 M, LDSTy, false, GlobalValue::InternalLinkage, PoisonValue::get(LDSTy),1242 VarName, nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS,1243 false);1244 SGV->setAlignment(StructAlign);1245 1246 DenseMap<GlobalVariable *, Constant *> Map;1247 Type *I32 = Type::getInt32Ty(Ctx);1248 for (size_t I = 0; I < LocalVars.size(); I++) {1249 GlobalVariable *GV = LocalVars[I];1250 Constant *GEPIdx[] = {ConstantInt::get(I32, 0), ConstantInt::get(I32, I)};1251 Constant *GEP = ConstantExpr::getGetElementPtr(LDSTy, SGV, GEPIdx, true);1252 if (IsPaddingField[I]) {1253 assert(GV->use_empty());1254 GV->eraseFromParent();1255 } else {1256 Map[GV] = GEP;1257 }1258 }1259 assert(Map.size() == LDSVarsToTransform.size());1260 return {SGV, std::move(Map)};1261 }1262 1263 template <typename PredicateTy>1264 static void replaceLDSVariablesWithStruct(1265 Module &M, DenseSet<GlobalVariable *> const &LDSVarsToTransformArg,1266 const LDSVariableReplacement &Replacement, PredicateTy Predicate) {1267 LLVMContext &Ctx = M.getContext();1268 const DataLayout &DL = M.getDataLayout();1269 1270 // A hack... we need to insert the aliasing info in a predictable order for1271 // lit tests. Would like to have them in a stable order already, ideally the1272 // same order they get allocated, which might mean an ordered set container1273 auto LDSVarsToTransform = sortByName(std::vector<GlobalVariable *>(1274 LDSVarsToTransformArg.begin(), LDSVarsToTransformArg.end()));1275 1276 // Create alias.scope and their lists. Each field in the new structure1277 // does not alias with all other fields.1278 SmallVector<MDNode *> AliasScopes;1279 SmallVector<Metadata *> NoAliasList;1280 const size_t NumberVars = LDSVarsToTransform.size();1281 if (NumberVars > 1) {1282 MDBuilder MDB(Ctx);1283 AliasScopes.reserve(NumberVars);1284 MDNode *Domain = MDB.createAnonymousAliasScopeDomain();1285 for (size_t I = 0; I < NumberVars; I++) {1286 MDNode *Scope = MDB.createAnonymousAliasScope(Domain);1287 AliasScopes.push_back(Scope);1288 }1289 NoAliasList.append(&AliasScopes[1], AliasScopes.end());1290 }1291 1292 // Replace uses of ith variable with a constantexpr to the corresponding1293 // field of the instance that will be allocated by AMDGPUMachineFunction1294 for (size_t I = 0; I < NumberVars; I++) {1295 GlobalVariable *GV = LDSVarsToTransform[I];1296 Constant *GEP = Replacement.LDSVarsToConstantGEP.at(GV);1297 1298 GV->replaceUsesWithIf(GEP, Predicate);1299 1300 APInt APOff(DL.getIndexTypeSizeInBits(GEP->getType()), 0);1301 GEP->stripAndAccumulateInBoundsConstantOffsets(DL, APOff);1302 uint64_t Offset = APOff.getZExtValue();1303 1304 Align A =1305 commonAlignment(Replacement.SGV->getAlign().valueOrOne(), Offset);1306 1307 if (I)1308 NoAliasList[I - 1] = AliasScopes[I - 1];1309 MDNode *NoAlias =1310 NoAliasList.empty() ? nullptr : MDNode::get(Ctx, NoAliasList);1311 MDNode *AliasScope =1312 AliasScopes.empty() ? nullptr : MDNode::get(Ctx, {AliasScopes[I]});1313 1314 refineUsesAlignmentAndAA(GEP, A, DL, AliasScope, NoAlias);1315 }1316 }1317 1318 static void refineUsesAlignmentAndAA(Value *Ptr, Align A,1319 const DataLayout &DL, MDNode *AliasScope,1320 MDNode *NoAlias, unsigned MaxDepth = 5) {1321 if (!MaxDepth || (A == 1 && !AliasScope))1322 return;1323 1324 ScopedNoAliasAAResult ScopedNoAlias;1325 1326 for (User *U : Ptr->users()) {1327 if (auto *I = dyn_cast<Instruction>(U)) {1328 if (AliasScope && I->mayReadOrWriteMemory()) {1329 MDNode *AS = I->getMetadata(LLVMContext::MD_alias_scope);1330 AS = (AS ? MDNode::getMostGenericAliasScope(AS, AliasScope)1331 : AliasScope);1332 I->setMetadata(LLVMContext::MD_alias_scope, AS);1333 1334 MDNode *NA = I->getMetadata(LLVMContext::MD_noalias);1335 1336 // Scoped aliases can originate from two different domains.1337 // First domain would be from LDS domain (created by this pass).1338 // All entries (LDS vars) into LDS struct will have same domain.1339 1340 // Second domain could be existing scoped aliases that are the1341 // results of noalias params and subsequent optimizations that1342 // may alter thesse sets.1343 1344 // We need to be careful how we create new alias sets, and1345 // have right scopes and domains for loads/stores of these new1346 // LDS variables. We intersect NoAlias set if alias sets belong1347 // to the same domain. This is the case if we have memcpy using1348 // LDS variables. Both src and dst of memcpy would belong to1349 // LDS struct, they donot alias.1350 // On the other hand, if one of the domains is LDS and other is1351 // existing domain prior to LDS, we need to have a union of all1352 // these aliases set to preserve existing aliasing information.1353 1354 SmallPtrSet<const MDNode *, 16> ExistingDomains, LDSDomains;1355 ScopedNoAlias.collectScopedDomains(NA, ExistingDomains);1356 ScopedNoAlias.collectScopedDomains(NoAlias, LDSDomains);1357 auto Intersection = set_intersection(ExistingDomains, LDSDomains);1358 if (Intersection.empty()) {1359 NA = NA ? MDNode::concatenate(NA, NoAlias) : NoAlias;1360 } else {1361 NA = NA ? MDNode::intersect(NA, NoAlias) : NoAlias;1362 }1363 I->setMetadata(LLVMContext::MD_noalias, NA);1364 }1365 }1366 1367 if (auto *LI = dyn_cast<LoadInst>(U)) {1368 LI->setAlignment(std::max(A, LI->getAlign()));1369 continue;1370 }1371 if (auto *SI = dyn_cast<StoreInst>(U)) {1372 if (SI->getPointerOperand() == Ptr)1373 SI->setAlignment(std::max(A, SI->getAlign()));1374 continue;1375 }1376 if (auto *AI = dyn_cast<AtomicRMWInst>(U)) {1377 // None of atomicrmw operations can work on pointers, but let's1378 // check it anyway in case it will or we will process ConstantExpr.1379 if (AI->getPointerOperand() == Ptr)1380 AI->setAlignment(std::max(A, AI->getAlign()));1381 continue;1382 }1383 if (auto *AI = dyn_cast<AtomicCmpXchgInst>(U)) {1384 if (AI->getPointerOperand() == Ptr)1385 AI->setAlignment(std::max(A, AI->getAlign()));1386 continue;1387 }1388 if (auto *GEP = dyn_cast<GetElementPtrInst>(U)) {1389 unsigned BitWidth = DL.getIndexTypeSizeInBits(GEP->getType());1390 APInt Off(BitWidth, 0);1391 if (GEP->getPointerOperand() == Ptr) {1392 Align GA;1393 if (GEP->accumulateConstantOffset(DL, Off))1394 GA = commonAlignment(A, Off.getLimitedValue());1395 refineUsesAlignmentAndAA(GEP, GA, DL, AliasScope, NoAlias,1396 MaxDepth - 1);1397 }1398 continue;1399 }1400 if (auto *I = dyn_cast<Instruction>(U)) {1401 if (I->getOpcode() == Instruction::BitCast ||1402 I->getOpcode() == Instruction::AddrSpaceCast)1403 refineUsesAlignmentAndAA(I, A, DL, AliasScope, NoAlias, MaxDepth - 1);1404 }1405 }1406 }1407};1408 1409class AMDGPULowerModuleLDSLegacy : public ModulePass {1410public:1411 const AMDGPUTargetMachine *TM;1412 static char ID;1413 1414 AMDGPULowerModuleLDSLegacy(const AMDGPUTargetMachine *TM = nullptr)1415 : ModulePass(ID), TM(TM) {}1416 1417 void getAnalysisUsage(AnalysisUsage &AU) const override {1418 if (!TM)1419 AU.addRequired<TargetPassConfig>();1420 }1421 1422 bool runOnModule(Module &M) override {1423 if (!TM) {1424 auto &TPC = getAnalysis<TargetPassConfig>();1425 TM = &TPC.getTM<AMDGPUTargetMachine>();1426 }1427 1428 return AMDGPULowerModuleLDS(*TM).runOnModule(M);1429 }1430};1431 1432} // namespace1433char AMDGPULowerModuleLDSLegacy::ID = 0;1434 1435char &llvm::AMDGPULowerModuleLDSLegacyPassID = AMDGPULowerModuleLDSLegacy::ID;1436 1437INITIALIZE_PASS_BEGIN(AMDGPULowerModuleLDSLegacy, DEBUG_TYPE,1438 "Lower uses of LDS variables from non-kernel functions",1439 false, false)1440INITIALIZE_PASS_DEPENDENCY(TargetPassConfig)1441INITIALIZE_PASS_END(AMDGPULowerModuleLDSLegacy, DEBUG_TYPE,1442 "Lower uses of LDS variables from non-kernel functions",1443 false, false)1444 1445ModulePass *1446llvm::createAMDGPULowerModuleLDSLegacyPass(const AMDGPUTargetMachine *TM) {1447 return new AMDGPULowerModuleLDSLegacy(TM);1448}1449 1450PreservedAnalyses AMDGPULowerModuleLDSPass::run(Module &M,1451 ModuleAnalysisManager &) {1452 return AMDGPULowerModuleLDS(TM).runOnModule(M) ? PreservedAnalyses::none()1453 : PreservedAnalyses::all();1454}1455