| //===-- AMDGPULowerModuleLDSPass.cpp ------------------------------*- C++ -*-=// |
| // |
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
| // See https://llvm.org/LICENSE.txt for license information. |
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
| // |
| //===----------------------------------------------------------------------===// |
| // |
| // This pass eliminates local data store, LDS, uses from non-kernel functions. |
| // LDS is contiguous memory allocated per kernel execution. |
| // |
| // Background. |
| // |
| // The programming model is global variables, or equivalently function local |
| // static variables, accessible from kernels or other functions. For uses from |
| // kernels this is straightforward - assign an integer to the kernel for the |
| // memory required by all the variables combined, allocate them within that. |
| // For uses from functions there are performance tradeoffs to choose between. |
| // |
| // This model means the GPU runtime can specify the amount of memory allocated. |
| // If this is more than the kernel assumed, the excess can be made available |
| // using a language specific feature, which IR represents as a variable with |
| // no initializer. This feature is not yet implemented for non-kernel functions. |
| // This lowering could be extended to handle that use case, but would probably |
| // require closer integration with promoteAllocaToLDS. |
| // |
| // Consequences of this GPU feature: |
| // - memory is limited and exceeding it halts compilation |
| // - a global accessed by one kernel exists independent of other kernels |
| // - a global exists independent of simultaneous execution of the same kernel |
| // - the address of the global may be different from different kernels as they |
| // do not alias, which permits only allocating variables they use |
| // - if the address is allowed to differ, functions need help to find it |
| // |
| // Uses from kernels are implemented here by grouping them in a per-kernel |
| // struct instance. This duplicates the variables, accurately modelling their |
| // aliasing properties relative to a single global representation. It also |
| // permits control over alignment via padding. |
| // |
| // Uses from functions are more complicated and the primary purpose of this |
| // IR pass. Several different lowering are chosen between to meet requirements |
| // to avoid allocating any LDS where it is not necessary, as that impacts |
| // occupancy and may fail the compilation, while not imposing overhead on a |
| // feature whose primary advantage over global memory is performance. The basic |
| // design goal is to avoid one kernel imposing overhead on another. |
| // |
| // Implementation. |
| // |
| // LDS variables with constant annotation or non-undef initializer are passed |
| // through unchanged for simplification or error diagnostics in later passes. |
| // Non-undef initializers are not yet implemented for LDS. |
| // |
| // LDS variables that are always allocated at the same address can be found |
| // by lookup at that address. Otherwise runtime information/cost is required. |
| // |
| // The simplest strategy possible is to group all LDS variables in a single |
| // struct and allocate that struct in every kernel such that the original |
| // variables are always at the same address. LDS is however a limited resource |
| // so this strategy is unusable in practice. It is not implemented here. |
| // |
| // Strategy | Precise allocation | Zero runtime cost | General purpose | |
| // --------+--------------------+-------------------+-----------------+ |
| // Module | No | Yes | Yes | |
| // Table | Yes | No | Yes | |
| // Kernel | Yes | Yes | No | |
| // Hybrid | Yes | Partial | Yes | |
| // |
| // Module spends LDS memory to save cycles. Table spends cycles and global |
| // memory to save LDS. Kernel is as fast as kernel allocation but only works |
| // for variables that are known reachable from a single kernel. Hybrid picks |
| // between all three. When forced to choose between LDS and cycles it minimises |
| // LDS use. |
| |
| // The "module" lowering implemented here finds LDS variables which are used by |
| // non-kernel functions and creates a new struct with a field for each of those |
| // LDS variables. Variables that are only used from kernels are excluded. |
| // Kernels that do not use this struct are annoteated with the attribute |
| // amdgpu-elide-module-lds which allows the back end to elide the allocation. |
| // |
| // The "table" lowering implemented here has three components. |
| // First kernels are assigned a unique integer identifier which is available in |
| // functions it calls through the intrinsic amdgcn_lds_kernel_id. The integer |
| // is passed through a specific SGPR, thus works with indirect calls. |
| // Second, each kernel allocates LDS variables independent of other kernels and |
| // writes the addresses it chose for each variable into an array in consistent |
| // order. If the kernel does not allocate a given variable, it writes undef to |
| // the corresponding array location. These arrays are written to a constant |
| // table in the order matching the kernel unique integer identifier. |
| // Third, uses from non-kernel functions are replaced with a table lookup using |
| // the intrinsic function to find the address of the variable. |
| // |
| // "Kernel" lowering is only applicable for variables that are unambiguously |
| // reachable from exactly one kernel. For those cases, accesses to the variable |
| // can be lowered to ConstantExpr address of a struct instance specific to that |
| // one kernel. This is zero cost in space and in compute. It will raise a fatal |
| // error on any variable that might be reachable from multiple kernels and is |
| // thus most easily used as part of the hybrid lowering strategy. |
| // |
| // Hybrid lowering is a mixture of the above. It uses the zero cost kernel |
| // lowering where it can. It lowers the variable accessed by the greatest |
| // number of kernels using the module strategy as that is free for the first |
| // variable. Any futher variables that can be lowered with the module strategy |
| // without incurring LDS memory overhead are. The remaining ones are lowered |
| // via table. |
| // |
| // Consequences |
| // - No heuristics or user controlled magic numbers, hybrid is the right choice |
| // - Kernels that don't use functions (or have had them all inlined) are not |
| // affected by any lowering for kernels that do. |
| // - Kernels that don't make indirect function calls are not affected by those |
| // that do. |
| // - Variables which are used by lots of kernels, e.g. those injected by a |
| // language runtime in most kernels, are expected to have no overhead |
| // - Implementations that instantiate templates per-kernel where those templates |
| // use LDS are expected to hit the "Kernel" lowering strategy |
| // - The runtime properties impose a cost in compiler implementation complexity |
| // |
| //===----------------------------------------------------------------------===// |
| |
| #include "AMDGPU.h" |
| #include "Utils/AMDGPUBaseInfo.h" |
| #include "Utils/AMDGPUMemoryUtils.h" |
| #include "llvm/ADT/BitVector.h" |
| #include "llvm/ADT/DenseMap.h" |
| #include "llvm/ADT/DenseSet.h" |
| #include "llvm/ADT/STLExtras.h" |
| #include "llvm/ADT/SetOperations.h" |
| #include "llvm/ADT/SetVector.h" |
| #include "llvm/Analysis/CallGraph.h" |
| #include "llvm/IR/Constants.h" |
| #include "llvm/IR/DerivedTypes.h" |
| #include "llvm/IR/IRBuilder.h" |
| #include "llvm/IR/InlineAsm.h" |
| #include "llvm/IR/Instructions.h" |
| #include "llvm/IR/IntrinsicsAMDGPU.h" |
| #include "llvm/IR/MDBuilder.h" |
| #include "llvm/InitializePasses.h" |
| #include "llvm/Pass.h" |
| #include "llvm/Support/CommandLine.h" |
| #include "llvm/Support/Debug.h" |
| #include "llvm/Support/OptimizedStructLayout.h" |
| #include "llvm/Transforms/Utils/BasicBlockUtils.h" |
| #include "llvm/Transforms/Utils/ModuleUtils.h" |
| |
| #include <tuple> |
| #include <vector> |
| |
| #include <cstdio> |
| |
| #define DEBUG_TYPE "amdgpu-lower-module-lds" |
| |
| using namespace llvm; |
| |
| namespace { |
| |
| cl::opt<bool> SuperAlignLDSGlobals( |
| "amdgpu-super-align-lds-globals", |
| cl::desc("Increase alignment of LDS if it is not on align boundary"), |
| cl::init(true), cl::Hidden); |
| |
| enum class LoweringKind { module, table, kernel, hybrid }; |
| cl::opt<LoweringKind> LoweringKindLoc( |
| "amdgpu-lower-module-lds-strategy", |
| cl::desc("Specify lowering strategy for function LDS access:"), cl::Hidden, |
| cl::init(LoweringKind::module), |
| cl::values( |
| clEnumValN(LoweringKind::table, "table", "Lower via table lookup"), |
| clEnumValN(LoweringKind::module, "module", "Lower via module struct"), |
| clEnumValN( |
| LoweringKind::kernel, "kernel", |
| "Lower variables reachable from one kernel, otherwise abort"), |
| clEnumValN(LoweringKind::hybrid, "hybrid", |
| "Lower via mixture of above strategies"))); |
| |
| bool isKernelLDS(const Function *F) { |
| // Some weirdness here. AMDGPU::isKernelCC does not call into |
| // AMDGPU::isKernel with the calling conv, it instead calls into |
| // isModuleEntryFunction which returns true for more calling conventions |
| // than AMDGPU::isKernel does. There's a FIXME on AMDGPU::isKernel. |
| // There's also a test that checks that the LDS lowering does not hit on |
| // a graphics shader, denoted amdgpu_ps, so stay with the limited case. |
| // Putting LDS in the name of the function to draw attention to this. |
| return AMDGPU::isKernel(F->getCallingConv()); |
| } |
| |
| class AMDGPULowerModuleLDS : public ModulePass { |
| |
| static void |
| removeLocalVarsFromUsedLists(Module &M, |
| const DenseSet<GlobalVariable *> &LocalVars) { |
| // The verifier rejects used lists containing an inttoptr of a constant |
| // so remove the variables from these lists before replaceAllUsesWith |
| SmallPtrSet<Constant *, 8> LocalVarsSet; |
| for (GlobalVariable *LocalVar : LocalVars) |
| LocalVarsSet.insert(cast<Constant>(LocalVar->stripPointerCasts())); |
| |
| removeFromUsedLists( |
| M, [&LocalVarsSet](Constant *C) { return LocalVarsSet.count(C); }); |
| |
| for (GlobalVariable *LocalVar : LocalVars) |
| LocalVar->removeDeadConstantUsers(); |
| } |
| |
| static void markUsedByKernel(IRBuilder<> &Builder, Function *Func, |
| GlobalVariable *SGV) { |
| // The llvm.amdgcn.module.lds instance is implicitly used by all kernels |
| // that might call a function which accesses a field within it. This is |
| // presently approximated to 'all kernels' if there are any such functions |
| // in the module. This implicit use is redefined as an explicit use here so |
| // that later passes, specifically PromoteAlloca, account for the required |
| // memory without any knowledge of this transform. |
| |
| // An operand bundle on llvm.donothing works because the call instruction |
| // survives until after the last pass that needs to account for LDS. It is |
| // better than inline asm as the latter survives until the end of codegen. A |
| // totally robust solution would be a function with the same semantics as |
| // llvm.donothing that takes a pointer to the instance and is lowered to a |
| // no-op after LDS is allocated, but that is not presently necessary. |
| |
| LLVMContext &Ctx = Func->getContext(); |
| |
| Builder.SetInsertPoint(Func->getEntryBlock().getFirstNonPHI()); |
| |
| FunctionType *FTy = FunctionType::get(Type::getVoidTy(Ctx), {}); |
| |
| Function *Decl = |
| Intrinsic::getDeclaration(Func->getParent(), Intrinsic::donothing, {}); |
| |
| Value *UseInstance[1] = {Builder.CreateInBoundsGEP( |
| SGV->getValueType(), SGV, ConstantInt::get(Type::getInt32Ty(Ctx), 0))}; |
| |
| Builder.CreateCall(FTy, Decl, {}, |
| {OperandBundleDefT<Value *>("ExplicitUse", UseInstance)}, |
| ""); |
| } |
| |
| static bool eliminateConstantExprUsesOfLDSFromAllInstructions(Module &M) { |
| // Constants are uniqued within LLVM. A ConstantExpr referring to a LDS |
| // global may have uses from multiple different functions as a result. |
| // This pass specialises LDS variables with respect to the kernel that |
| // allocates them. |
| |
| // This is semantically equivalent to: |
| // for (auto &F : M.functions()) |
| // for (auto &BB : F) |
| // for (auto &I : BB) |
| // for (Use &Op : I.operands()) |
| // if (constantExprUsesLDS(Op)) |
| // replaceConstantExprInFunction(I, Op); |
| |
| bool Changed = false; |
| |
| // Find all ConstantExpr that are direct users of an LDS global |
| SmallVector<ConstantExpr *> Stack; |
| for (auto &GV : M.globals()) |
| if (AMDGPU::isLDSVariableToLower(GV)) |
| for (User *U : GV.users()) |
| if (ConstantExpr *C = dyn_cast<ConstantExpr>(U)) |
| Stack.push_back(C); |
| |
| // Expand to include constexpr users of direct users |
| SetVector<ConstantExpr *> ConstExprUsersOfLDS; |
| while (!Stack.empty()) { |
| ConstantExpr *V = Stack.pop_back_val(); |
| if (ConstExprUsersOfLDS.contains(V)) |
| continue; |
| |
| ConstExprUsersOfLDS.insert(V); |
| |
| for (auto *Nested : V->users()) |
| if (ConstantExpr *CE = dyn_cast<ConstantExpr>(Nested)) |
| Stack.push_back(CE); |
| } |
| |
| // Find all instructions that use any of the ConstExpr users of LDS |
| SetVector<Instruction *> InstructionWorklist; |
| for (ConstantExpr *CE : ConstExprUsersOfLDS) |
| for (User *U : CE->users()) |
| if (auto *I = dyn_cast<Instruction>(U)) |
| InstructionWorklist.insert(I); |
| |
| // Replace those ConstExpr operands with instructions |
| while (!InstructionWorklist.empty()) { |
| Instruction *I = InstructionWorklist.pop_back_val(); |
| for (Use &U : I->operands()) { |
| |
| auto *BI = I; |
| if (auto *Phi = dyn_cast<PHINode>(I)) { |
| BasicBlock *BB = Phi->getIncomingBlock(U); |
| BasicBlock::iterator It = BB->getFirstInsertionPt(); |
| assert(It != BB->end() && "Unexpected empty basic block"); |
| BI = &(*(It)); |
| } |
| |
| if (ConstantExpr *C = dyn_cast<ConstantExpr>(U.get())) { |
| if (ConstExprUsersOfLDS.contains(C)) { |
| Changed = true; |
| Instruction *NI = C->getAsInstruction(BI); |
| InstructionWorklist.insert(NI); |
| U.set(NI); |
| C->removeDeadConstantUsers(); |
| } |
| } |
| } |
| } |
| |
| return Changed; |
| } |
| |
| public: |
| static char ID; |
| |
| AMDGPULowerModuleLDS() : ModulePass(ID) { |
| initializeAMDGPULowerModuleLDSPass(*PassRegistry::getPassRegistry()); |
| } |
| |
| using FunctionVariableMap = DenseMap<Function *, DenseSet<GlobalVariable *>>; |
| |
| using VariableFunctionMap = DenseMap<GlobalVariable *, DenseSet<Function *>>; |
| |
| static void getUsesOfLDSByFunction(CallGraph const &CG, Module &M, |
| FunctionVariableMap &kernels, |
| FunctionVariableMap &functions) { |
| |
| // Get uses from the current function, excluding uses by called functions |
| // Two output variables to avoid walking the globals list twice |
| for (auto &GV : M.globals()) { |
| if (!AMDGPU::isLDSVariableToLower(GV)) { |
| continue; |
| } |
| |
| SmallVector<User *, 16> Stack(GV.users()); |
| for (User *V : GV.users()) { |
| if (auto *I = dyn_cast<Instruction>(V)) { |
| Function *F = I->getFunction(); |
| if (isKernelLDS(F)) { |
| kernels[F].insert(&GV); |
| } else { |
| functions[F].insert(&GV); |
| } |
| } |
| } |
| } |
| } |
| |
| struct LDSUsesInfoTy { |
| FunctionVariableMap direct_access; |
| FunctionVariableMap indirect_access; |
| }; |
| |
| static LDSUsesInfoTy getTransitiveUsesOfLDS(CallGraph const &CG, Module &M) { |
| |
| FunctionVariableMap direct_map_kernel; |
| FunctionVariableMap direct_map_function; |
| getUsesOfLDSByFunction(CG, M, direct_map_kernel, direct_map_function); |
| |
| // Collect variables that are used by functions whose address has escaped |
| DenseSet<GlobalVariable *> VariablesReachableThroughFunctionPointer; |
| for (Function &F : M.functions()) { |
| if (!isKernelLDS(&F)) |
| if (F.hasAddressTaken(nullptr, |
| /* IgnoreCallbackUses */ false, |
| /* IgnoreAssumeLikeCalls */ false, |
| /* IgnoreLLVMUsed */ true, |
| /* IgnoreArcAttachedCall */ false)) { |
| set_union(VariablesReachableThroughFunctionPointer, |
| direct_map_function[&F]); |
| } |
| } |
| |
| auto functionMakesUnknownCall = [&](const Function *F) -> bool { |
| assert(!F->isDeclaration()); |
| for (CallGraphNode::CallRecord R : *CG[F]) { |
| if (!R.second->getFunction()) { |
| return true; |
| } |
| } |
| return false; |
| }; |
| |
| // Work out which variables are reachable through function calls |
| FunctionVariableMap transitive_map_function = direct_map_function; |
| |
| // If the function makes any unknown call, assume the worst case that it can |
| // access all variables accessed by functions whose address escaped |
| for (Function &F : M.functions()) { |
| if (!F.isDeclaration() && functionMakesUnknownCall(&F)) { |
| if (!isKernelLDS(&F)) { |
| set_union(transitive_map_function[&F], |
| VariablesReachableThroughFunctionPointer); |
| } |
| } |
| } |
| |
| // Direct implementation of collecting all variables reachable from each |
| // function |
| for (Function &Func : M.functions()) { |
| if (Func.isDeclaration() || isKernelLDS(&Func)) |
| continue; |
| |
| DenseSet<Function *> seen; // catches cycles |
| SmallVector<Function *, 4> wip{&Func}; |
| |
| while (!wip.empty()) { |
| Function *F = wip.pop_back_val(); |
| |
| // Can accelerate this by referring to transitive map for functions that |
| // have already been computed, with more care than this |
| set_union(transitive_map_function[&Func], direct_map_function[F]); |
| |
| for (CallGraphNode::CallRecord R : *CG[F]) { |
| Function *ith = R.second->getFunction(); |
| if (ith) { |
| if (!seen.contains(ith)) { |
| seen.insert(ith); |
| wip.push_back(ith); |
| } |
| } |
| } |
| } |
| } |
| |
| // direct_map_kernel lists which variables are used by the kernel |
| // find the variables which are used through a function call |
| FunctionVariableMap indirect_map_kernel; |
| |
| for (Function &Func : M.functions()) { |
| if (Func.isDeclaration() || !isKernelLDS(&Func)) |
| continue; |
| |
| for (CallGraphNode::CallRecord R : *CG[&Func]) { |
| Function *ith = R.second->getFunction(); |
| if (ith) { |
| set_union(indirect_map_kernel[&Func], transitive_map_function[ith]); |
| } else { |
| set_union(indirect_map_kernel[&Func], |
| VariablesReachableThroughFunctionPointer); |
| } |
| } |
| } |
| |
| return {std::move(direct_map_kernel), std::move(indirect_map_kernel)}; |
| } |
| |
| struct LDSVariableReplacement { |
| GlobalVariable *SGV = nullptr; |
| DenseMap<GlobalVariable *, Constant *> LDSVarsToConstantGEP; |
| }; |
| |
| // remap from lds global to a constantexpr gep to where it has been moved to |
| // for each kernel |
| // an array with an element for each kernel containing where the corresponding |
| // variable was remapped to |
| |
| static Constant *getAddressesOfVariablesInKernel( |
| LLVMContext &Ctx, ArrayRef<GlobalVariable *> Variables, |
| DenseMap<GlobalVariable *, Constant *> &LDSVarsToConstantGEP) { |
| // Create a ConstantArray containing the address of each Variable within the |
| // kernel corresponding to LDSVarsToConstantGEP, or poison if that kernel |
| // does not allocate it |
| // TODO: Drop the ptrtoint conversion |
| |
| Type *I32 = Type::getInt32Ty(Ctx); |
| |
| ArrayType *KernelOffsetsType = ArrayType::get(I32, Variables.size()); |
| |
| SmallVector<Constant *> Elements; |
| for (size_t i = 0; i < Variables.size(); i++) { |
| GlobalVariable *GV = Variables[i]; |
| if (LDSVarsToConstantGEP.count(GV) != 0) { |
| auto elt = ConstantExpr::getPtrToInt(LDSVarsToConstantGEP[GV], I32); |
| Elements.push_back(elt); |
| } else { |
| Elements.push_back(PoisonValue::get(I32)); |
| } |
| } |
| return ConstantArray::get(KernelOffsetsType, Elements); |
| } |
| |
| static GlobalVariable *buildLookupTable( |
| Module &M, ArrayRef<GlobalVariable *> Variables, |
| ArrayRef<Function *> kernels, |
| DenseMap<Function *, LDSVariableReplacement> &KernelToReplacement) { |
| if (Variables.empty()) { |
| return nullptr; |
| } |
| LLVMContext &Ctx = M.getContext(); |
| |
| const size_t NumberVariables = Variables.size(); |
| const size_t NumberKernels = kernels.size(); |
| |
| ArrayType *KernelOffsetsType = |
| ArrayType::get(Type::getInt32Ty(Ctx), NumberVariables); |
| |
| ArrayType *AllKernelsOffsetsType = |
| ArrayType::get(KernelOffsetsType, NumberKernels); |
| |
| std::vector<Constant *> overallConstantExprElts(NumberKernels); |
| for (size_t i = 0; i < NumberKernels; i++) { |
| LDSVariableReplacement Replacement = KernelToReplacement[kernels[i]]; |
| overallConstantExprElts[i] = getAddressesOfVariablesInKernel( |
| Ctx, Variables, Replacement.LDSVarsToConstantGEP); |
| } |
| |
| Constant *init = |
| ConstantArray::get(AllKernelsOffsetsType, overallConstantExprElts); |
| |
| return new GlobalVariable( |
| M, AllKernelsOffsetsType, true, GlobalValue::InternalLinkage, init, |
| "llvm.amdgcn.lds.offset.table", nullptr, GlobalValue::NotThreadLocal, |
| AMDGPUAS::CONSTANT_ADDRESS); |
| } |
| |
| void replaceUsesInInstructionsWithTableLookup( |
| Module &M, ArrayRef<GlobalVariable *> ModuleScopeVariables, |
| GlobalVariable *LookupTable) { |
| |
| LLVMContext &Ctx = M.getContext(); |
| IRBuilder<> Builder(Ctx); |
| Type *I32 = Type::getInt32Ty(Ctx); |
| |
| // Accesses from a function use the amdgcn_lds_kernel_id intrinsic which |
| // lowers to a read from a live in register. Emit it once in the entry |
| // block to spare deduplicating it later. |
| |
| DenseMap<Function *, Value *> tableKernelIndexCache; |
| auto getTableKernelIndex = [&](Function *F) -> Value * { |
| if (tableKernelIndexCache.count(F) == 0) { |
| LLVMContext &Ctx = M.getContext(); |
| FunctionType *FTy = FunctionType::get(Type::getInt32Ty(Ctx), {}); |
| Function *Decl = |
| Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_lds_kernel_id, {}); |
| |
| BasicBlock::iterator it = |
| F->getEntryBlock().getFirstNonPHIOrDbgOrAlloca(); |
| Instruction &i = *it; |
| Builder.SetInsertPoint(&i); |
| |
| tableKernelIndexCache[F] = Builder.CreateCall(FTy, Decl, {}); |
| } |
| |
| return tableKernelIndexCache[F]; |
| }; |
| |
| for (size_t Index = 0; Index < ModuleScopeVariables.size(); Index++) { |
| auto *GV = ModuleScopeVariables[Index]; |
| |
| for (Use &U : make_early_inc_range(GV->uses())) { |
| auto *I = dyn_cast<Instruction>(U.getUser()); |
| if (!I) |
| continue; |
| |
| Value *tableKernelIndex = getTableKernelIndex(I->getFunction()); |
| |
| // So if the phi uses this value multiple times, what does this look |
| // like? |
| if (auto *Phi = dyn_cast<PHINode>(I)) { |
| BasicBlock *BB = Phi->getIncomingBlock(U); |
| Builder.SetInsertPoint(&(*(BB->getFirstInsertionPt()))); |
| } else { |
| Builder.SetInsertPoint(I); |
| } |
| |
| Value *GEPIdx[3] = { |
| ConstantInt::get(I32, 0), |
| tableKernelIndex, |
| ConstantInt::get(I32, Index), |
| }; |
| |
| Value *Address = Builder.CreateInBoundsGEP( |
| LookupTable->getValueType(), LookupTable, GEPIdx, GV->getName()); |
| |
| Value *loaded = Builder.CreateLoad(I32, Address); |
| |
| Value *replacement = |
| Builder.CreateIntToPtr(loaded, GV->getType(), GV->getName()); |
| |
| U.set(replacement); |
| } |
| } |
| } |
| |
| static DenseSet<Function *> kernelsThatIndirectlyAccessAnyOfPassedVariables( |
| Module &M, LDSUsesInfoTy &LDSUsesInfo, |
| DenseSet<GlobalVariable *> const &VariableSet) { |
| |
| DenseSet<Function *> KernelSet; |
| |
| if (VariableSet.empty()) return KernelSet; |
| |
| for (Function &Func : M.functions()) { |
| if (Func.isDeclaration() || !isKernelLDS(&Func)) |
| continue; |
| for (GlobalVariable *GV : LDSUsesInfo.indirect_access[&Func]) { |
| if (VariableSet.contains(GV)) { |
| KernelSet.insert(&Func); |
| break; |
| } |
| } |
| } |
| |
| return KernelSet; |
| } |
| |
| static GlobalVariable * |
| chooseBestVariableForModuleStrategy(const DataLayout &DL, |
| VariableFunctionMap &LDSVars) { |
| // Find the global variable with the most indirect uses from kernels |
| |
| struct CandidateTy { |
| GlobalVariable *GV = nullptr; |
| size_t UserCount = 0; |
| size_t Size = 0; |
| |
| CandidateTy() = default; |
| |
| CandidateTy(GlobalVariable *GV, uint64_t UserCount, uint64_t AllocSize) |
| : GV(GV), UserCount(UserCount), Size(AllocSize) {} |
| |
| bool operator<(const CandidateTy &Other) const { |
| // Fewer users makes module scope variable less attractive |
| if (UserCount < Other.UserCount) { |
| return true; |
| } |
| if (UserCount > Other.UserCount) { |
| return false; |
| } |
| |
| // Bigger makes module scope variable less attractive |
| if (Size < Other.Size) { |
| return false; |
| } |
| |
| if (Size > Other.Size) { |
| return true; |
| } |
| |
| // Arbitrary but consistent |
| return GV->getName() < Other.GV->getName(); |
| } |
| }; |
| |
| CandidateTy MostUsed; |
| |
| for (auto &K : LDSVars) { |
| GlobalVariable *GV = K.first; |
| if (K.second.size() <= 1) { |
| // A variable reachable by only one kernel is best lowered with kernel |
| // strategy |
| continue; |
| } |
| CandidateTy Candidate(GV, K.second.size(), |
| DL.getTypeAllocSize(GV->getValueType()).getFixedValue()); |
| if (MostUsed < Candidate) |
| MostUsed = Candidate; |
| } |
| |
| return MostUsed.GV; |
| } |
| |
| bool runOnModule(Module &M) override { |
| LLVMContext &Ctx = M.getContext(); |
| CallGraph CG = CallGraph(M); |
| bool Changed = superAlignLDSGlobals(M); |
| |
| Changed |= eliminateConstantExprUsesOfLDSFromAllInstructions(M); |
| |
| Changed = true; // todo: narrow this down |
| |
| // For each kernel, what variables does it access directly or through |
| // callees |
| LDSUsesInfoTy LDSUsesInfo = getTransitiveUsesOfLDS(CG, M); |
| |
| // For each variable accessed through callees, which kernels access it |
| VariableFunctionMap LDSToKernelsThatNeedToAccessItIndirectly; |
| for (auto &K : LDSUsesInfo.indirect_access) { |
| Function *F = K.first; |
| assert(isKernelLDS(F)); |
| for (GlobalVariable *GV : K.second) { |
| LDSToKernelsThatNeedToAccessItIndirectly[GV].insert(F); |
| } |
| } |
| |
| // Partition variables into the different strategies |
| DenseSet<GlobalVariable *> ModuleScopeVariables; |
| DenseSet<GlobalVariable *> TableLookupVariables; |
| DenseSet<GlobalVariable *> KernelAccessVariables; |
| |
| { |
| GlobalVariable *HybridModuleRoot = |
| LoweringKindLoc != LoweringKind::hybrid |
| ? nullptr |
| : chooseBestVariableForModuleStrategy( |
| M.getDataLayout(), |
| LDSToKernelsThatNeedToAccessItIndirectly); |
| |
| DenseSet<Function *> const EmptySet; |
| DenseSet<Function *> const &HybridModuleRootKernels = |
| HybridModuleRoot |
| ? LDSToKernelsThatNeedToAccessItIndirectly[HybridModuleRoot] |
| : EmptySet; |
| |
| for (auto &K : LDSToKernelsThatNeedToAccessItIndirectly) { |
| // Each iteration of this loop assigns exactly one global variable to |
| // exactly one of the implementation strategies. |
| |
| GlobalVariable *GV = K.first; |
| assert(AMDGPU::isLDSVariableToLower(*GV)); |
| assert(K.second.size() != 0); |
| |
| switch (LoweringKindLoc) { |
| case LoweringKind::module: |
| ModuleScopeVariables.insert(GV); |
| break; |
| |
| case LoweringKind::table: |
| TableLookupVariables.insert(GV); |
| break; |
| |
| case LoweringKind::kernel: |
| if (K.second.size() == 1) { |
| KernelAccessVariables.insert(GV); |
| } else { |
| report_fatal_error( |
| "cannot lower LDS '" + GV->getName() + |
| "' to kernel access as it is reachable from multiple kernels"); |
| } |
| break; |
| |
| case LoweringKind::hybrid: { |
| if (GV == HybridModuleRoot) { |
| assert(K.second.size() != 1); |
| ModuleScopeVariables.insert(GV); |
| } else if (K.second.size() == 1) { |
| KernelAccessVariables.insert(GV); |
| } else if (set_is_subset(K.second, HybridModuleRootKernels)) { |
| ModuleScopeVariables.insert(GV); |
| } else { |
| TableLookupVariables.insert(GV); |
| } |
| break; |
| } |
| } |
| } |
| |
| assert(ModuleScopeVariables.size() + TableLookupVariables.size() + |
| KernelAccessVariables.size() == |
| LDSToKernelsThatNeedToAccessItIndirectly.size()); |
| } // Variables have now been partitioned into the three lowering strategies. |
| |
| // If the kernel accesses a variable that is going to be stored in the |
| // module instance through a call then that kernel needs to allocate the |
| // module instance |
| DenseSet<Function *> KernelsThatAllocateModuleLDS = |
| kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo, |
| ModuleScopeVariables); |
| DenseSet<Function *> KernelsThatAllocateTableLDS = |
| kernelsThatIndirectlyAccessAnyOfPassedVariables(M, LDSUsesInfo, |
| TableLookupVariables); |
| |
| if (!ModuleScopeVariables.empty()) { |
| LDSVariableReplacement ModuleScopeReplacement = |
| createLDSVariableReplacement(M, "llvm.amdgcn.module.lds", |
| ModuleScopeVariables); |
| |
| appendToCompilerUsed(M, |
| {static_cast<GlobalValue *>( |
| ConstantExpr::getPointerBitCastOrAddrSpaceCast( |
| cast<Constant>(ModuleScopeReplacement.SGV), |
| Type::getInt8PtrTy(Ctx)))}); |
| |
| // historic |
| removeLocalVarsFromUsedLists(M, ModuleScopeVariables); |
| |
| // Replace all uses of module scope variable from non-kernel functions |
| replaceLDSVariablesWithStruct( |
| M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) { |
| Instruction *I = dyn_cast<Instruction>(U.getUser()); |
| if (!I) { |
| return false; |
| } |
| Function *F = I->getFunction(); |
| return !isKernelLDS(F); |
| }); |
| |
| // Replace uses of module scope variable from kernel functions that |
| // allocate the module scope variable, otherwise leave them unchanged |
| // Record on each kernel whether the module scope global is used by it |
| |
| LLVMContext &Ctx = M.getContext(); |
| IRBuilder<> Builder(Ctx); |
| |
| for (Function &Func : M.functions()) { |
| if (Func.isDeclaration() || !isKernelLDS(&Func)) |
| continue; |
| |
| if (KernelsThatAllocateModuleLDS.contains(&Func)) { |
| replaceLDSVariablesWithStruct( |
| M, ModuleScopeVariables, ModuleScopeReplacement, [&](Use &U) { |
| Instruction *I = dyn_cast<Instruction>(U.getUser()); |
| if (!I) { |
| return false; |
| } |
| Function *F = I->getFunction(); |
| return F == &Func; |
| }); |
| |
| markUsedByKernel(Builder, &Func, ModuleScopeReplacement.SGV); |
| |
| } else { |
| Func.addFnAttr("amdgpu-elide-module-lds"); |
| } |
| } |
| } |
| |
| // Create a struct for each kernel for the non-module-scope variables |
| DenseMap<Function *, LDSVariableReplacement> KernelToReplacement; |
| for (Function &Func : M.functions()) { |
| if (Func.isDeclaration() || !isKernelLDS(&Func)) |
| continue; |
| |
| DenseSet<GlobalVariable *> KernelUsedVariables; |
| for (auto &v : LDSUsesInfo.direct_access[&Func]) { |
| KernelUsedVariables.insert(v); |
| } |
| for (auto &v : LDSUsesInfo.indirect_access[&Func]) { |
| KernelUsedVariables.insert(v); |
| } |
| |
| // Variables allocated in module lds must all resolve to that struct, |
| // not to the per-kernel instance. |
| if (KernelsThatAllocateModuleLDS.contains(&Func)) { |
| for (GlobalVariable *v : ModuleScopeVariables) { |
| KernelUsedVariables.erase(v); |
| } |
| } |
| |
| if (KernelUsedVariables.empty()) { |
| // Either used no LDS, or all the LDS it used was also in module |
| continue; |
| } |
| |
| // The association between kernel function and LDS struct is done by |
| // symbol name, which only works if the function in question has a |
| // name This is not expected to be a problem in practice as kernels |
| // are called by name making anonymous ones (which are named by the |
| // backend) difficult to use. This does mean that llvm test cases need |
| // to name the kernels. |
| if (!Func.hasName()) { |
| report_fatal_error("Anonymous kernels cannot use LDS variables"); |
| } |
| |
| std::string VarName = |
| (Twine("llvm.amdgcn.kernel.") + Func.getName() + ".lds").str(); |
| |
| auto Replacement = |
| createLDSVariableReplacement(M, VarName, KernelUsedVariables); |
| |
| // remove preserves existing codegen |
| removeLocalVarsFromUsedLists(M, KernelUsedVariables); |
| KernelToReplacement[&Func] = Replacement; |
| |
| // Rewrite uses within kernel to the new struct |
| replaceLDSVariablesWithStruct( |
| M, KernelUsedVariables, Replacement, [&Func](Use &U) { |
| Instruction *I = dyn_cast<Instruction>(U.getUser()); |
| return I && I->getFunction() == &Func; |
| }); |
| } |
| |
| // Lower zero cost accesses to the kernel instances just created |
| for (auto &GV : KernelAccessVariables) { |
| auto &funcs = LDSToKernelsThatNeedToAccessItIndirectly[GV]; |
| assert(funcs.size() == 1); // Only one kernel can access it |
| LDSVariableReplacement Replacement = |
| KernelToReplacement[*(funcs.begin())]; |
| |
| DenseSet<GlobalVariable *> Vec; |
| Vec.insert(GV); |
| |
| replaceLDSVariablesWithStruct(M, Vec, Replacement, [](Use &U) { |
| return isa<Instruction>(U.getUser()); |
| }); |
| } |
| |
| if (!KernelsThatAllocateTableLDS.empty()) { |
| // Collect the kernels that allocate table lookup LDS |
| std::vector<Function *> OrderedKernels; |
| { |
| for (Function &Func : M.functions()) { |
| if (Func.isDeclaration()) |
| continue; |
| if (!isKernelLDS(&Func)) |
| continue; |
| |
| if (KernelsThatAllocateTableLDS.contains(&Func)) { |
| assert(Func.hasName()); // else fatal error earlier |
| OrderedKernels.push_back(&Func); |
| } |
| } |
| |
| // Put them in an arbitrary but reproducible order |
| llvm::sort(OrderedKernels.begin(), OrderedKernels.end(), |
| [](const Function *lhs, const Function *rhs) -> bool { |
| return lhs->getName() < rhs->getName(); |
| }); |
| |
| // Annotate the kernels with their order in this vector |
| LLVMContext &Ctx = M.getContext(); |
| IRBuilder<> Builder(Ctx); |
| |
| if (OrderedKernels.size() > UINT32_MAX) { |
| // 32 bit keeps it in one SGPR. > 2**32 kernels won't fit on the GPU |
| report_fatal_error("Unimplemented LDS lowering for > 2**32 kernels"); |
| } |
| |
| for (size_t i = 0; i < OrderedKernels.size(); i++) { |
| Metadata *AttrMDArgs[1] = { |
| ConstantAsMetadata::get(Builder.getInt32(i)), |
| }; |
| OrderedKernels[i]->setMetadata("llvm.amdgcn.lds.kernel.id", |
| MDNode::get(Ctx, AttrMDArgs)); |
| |
| markUsedByKernel(Builder, OrderedKernels[i], |
| KernelToReplacement[OrderedKernels[i]].SGV); |
| } |
| } |
| |
| // The order must be consistent between lookup table and accesses to |
| // lookup table |
| std::vector<GlobalVariable *> TableLookupVariablesOrdered( |
| TableLookupVariables.begin(), TableLookupVariables.end()); |
| llvm::sort(TableLookupVariablesOrdered.begin(), |
| TableLookupVariablesOrdered.end(), |
| [](const GlobalVariable *lhs, const GlobalVariable *rhs) { |
| return lhs->getName() < rhs->getName(); |
| }); |
| |
| GlobalVariable *LookupTable = buildLookupTable( |
| M, TableLookupVariablesOrdered, OrderedKernels, KernelToReplacement); |
| replaceUsesInInstructionsWithTableLookup(M, TableLookupVariablesOrdered, |
| LookupTable); |
| } |
| |
| for (auto &GV : make_early_inc_range(M.globals())) |
| if (AMDGPU::isLDSVariableToLower(GV)) { |
| |
| // probably want to remove from used lists |
| GV.removeDeadConstantUsers(); |
| if (GV.use_empty()) |
| GV.eraseFromParent(); |
| } |
| |
| return Changed; |
| } |
| |
| private: |
| // Increase the alignment of LDS globals if necessary to maximise the chance |
| // that we can use aligned LDS instructions to access them. |
| static bool superAlignLDSGlobals(Module &M) { |
| const DataLayout &DL = M.getDataLayout(); |
| bool Changed = false; |
| if (!SuperAlignLDSGlobals) { |
| return Changed; |
| } |
| |
| for (auto &GV : M.globals()) { |
| if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) { |
| // Only changing alignment of LDS variables |
| continue; |
| } |
| if (!GV.hasInitializer()) { |
| // cuda/hip extern __shared__ variable, leave alignment alone |
| continue; |
| } |
| |
| Align Alignment = AMDGPU::getAlign(DL, &GV); |
| TypeSize GVSize = DL.getTypeAllocSize(GV.getValueType()); |
| |
| if (GVSize > 8) { |
| // We might want to use a b96 or b128 load/store |
| Alignment = std::max(Alignment, Align(16)); |
| } else if (GVSize > 4) { |
| // We might want to use a b64 load/store |
| Alignment = std::max(Alignment, Align(8)); |
| } else if (GVSize > 2) { |
| // We might want to use a b32 load/store |
| Alignment = std::max(Alignment, Align(4)); |
| } else if (GVSize > 1) { |
| // We might want to use a b16 load/store |
| Alignment = std::max(Alignment, Align(2)); |
| } |
| |
| if (Alignment != AMDGPU::getAlign(DL, &GV)) { |
| Changed = true; |
| GV.setAlignment(Alignment); |
| } |
| } |
| return Changed; |
| } |
| |
| static LDSVariableReplacement createLDSVariableReplacement( |
| Module &M, std::string VarName, |
| DenseSet<GlobalVariable *> const &LDSVarsToTransform) { |
| // Create a struct instance containing LDSVarsToTransform and map from those |
| // variables to ConstantExprGEP |
| // Variables may be introduced to meet alignment requirements. No aliasing |
| // metadata is useful for these as they have no uses. Erased before return. |
| |
| LLVMContext &Ctx = M.getContext(); |
| const DataLayout &DL = M.getDataLayout(); |
| assert(!LDSVarsToTransform.empty()); |
| |
| SmallVector<OptimizedStructLayoutField, 8> LayoutFields; |
| LayoutFields.reserve(LDSVarsToTransform.size()); |
| { |
| // The order of fields in this struct depends on the order of |
| // varables in the argument which varies when changing how they |
| // are identified, leading to spurious test breakage. |
| std::vector<GlobalVariable *> Sorted(LDSVarsToTransform.begin(), |
| LDSVarsToTransform.end()); |
| llvm::sort(Sorted.begin(), Sorted.end(), |
| [](const GlobalVariable *lhs, const GlobalVariable *rhs) { |
| return lhs->getName() < rhs->getName(); |
| }); |
| for (GlobalVariable *GV : Sorted) { |
| OptimizedStructLayoutField F(GV, |
| DL.getTypeAllocSize(GV->getValueType()), |
| AMDGPU::getAlign(DL, GV)); |
| LayoutFields.emplace_back(F); |
| } |
| } |
| |
| performOptimizedStructLayout(LayoutFields); |
| |
| std::vector<GlobalVariable *> LocalVars; |
| BitVector IsPaddingField; |
| LocalVars.reserve(LDSVarsToTransform.size()); // will be at least this large |
| IsPaddingField.reserve(LDSVarsToTransform.size()); |
| { |
| uint64_t CurrentOffset = 0; |
| for (size_t I = 0; I < LayoutFields.size(); I++) { |
| GlobalVariable *FGV = static_cast<GlobalVariable *>( |
| const_cast<void *>(LayoutFields[I].Id)); |
| Align DataAlign = LayoutFields[I].Alignment; |
| |
| uint64_t DataAlignV = DataAlign.value(); |
| if (uint64_t Rem = CurrentOffset % DataAlignV) { |
| uint64_t Padding = DataAlignV - Rem; |
| |
| // Append an array of padding bytes to meet alignment requested |
| // Note (o + (a - (o % a)) ) % a == 0 |
| // (offset + Padding ) % align == 0 |
| |
| Type *ATy = ArrayType::get(Type::getInt8Ty(Ctx), Padding); |
| LocalVars.push_back(new GlobalVariable( |
| M, ATy, false, GlobalValue::InternalLinkage, UndefValue::get(ATy), |
| "", nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS, |
| false)); |
| IsPaddingField.push_back(true); |
| CurrentOffset += Padding; |
| } |
| |
| LocalVars.push_back(FGV); |
| IsPaddingField.push_back(false); |
| CurrentOffset += LayoutFields[I].Size; |
| } |
| } |
| |
| std::vector<Type *> LocalVarTypes; |
| LocalVarTypes.reserve(LocalVars.size()); |
| std::transform( |
| LocalVars.cbegin(), LocalVars.cend(), std::back_inserter(LocalVarTypes), |
| [](const GlobalVariable *V) -> Type * { return V->getValueType(); }); |
| |
| StructType *LDSTy = StructType::create(Ctx, LocalVarTypes, VarName + ".t"); |
| |
| Align StructAlign = AMDGPU::getAlign(DL, LocalVars[0]); |
| |
| GlobalVariable *SGV = new GlobalVariable( |
| M, LDSTy, false, GlobalValue::InternalLinkage, UndefValue::get(LDSTy), |
| VarName, nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS, |
| false); |
| SGV->setAlignment(StructAlign); |
| |
| DenseMap<GlobalVariable *, Constant *> Map; |
| Type *I32 = Type::getInt32Ty(Ctx); |
| for (size_t I = 0; I < LocalVars.size(); I++) { |
| GlobalVariable *GV = LocalVars[I]; |
| Constant *GEPIdx[] = {ConstantInt::get(I32, 0), ConstantInt::get(I32, I)}; |
| Constant *GEP = ConstantExpr::getGetElementPtr(LDSTy, SGV, GEPIdx, true); |
| if (IsPaddingField[I]) { |
| assert(GV->use_empty()); |
| GV->eraseFromParent(); |
| } else { |
| Map[GV] = GEP; |
| } |
| } |
| assert(Map.size() == LDSVarsToTransform.size()); |
| return {SGV, std::move(Map)}; |
| } |
| |
| template <typename PredicateTy> |
| void replaceLDSVariablesWithStruct( |
| Module &M, DenseSet<GlobalVariable *> const &LDSVarsToTransformArg, |
| LDSVariableReplacement Replacement, PredicateTy Predicate) { |
| LLVMContext &Ctx = M.getContext(); |
| const DataLayout &DL = M.getDataLayout(); |
| |
| // A hack... we need to insert the aliasing info in a predictable order for |
| // lit tests. Would like to have them in a stable order already, ideally the |
| // same order they get allocated, which might mean an ordered set container |
| std::vector<GlobalVariable *> LDSVarsToTransform( |
| LDSVarsToTransformArg.begin(), LDSVarsToTransformArg.end()); |
| llvm::sort(LDSVarsToTransform.begin(), LDSVarsToTransform.end(), |
| [](const GlobalVariable *lhs, const GlobalVariable *rhs) { |
| return lhs->getName() < rhs->getName(); |
| }); |
| |
| // Create alias.scope and their lists. Each field in the new structure |
| // does not alias with all other fields. |
| SmallVector<MDNode *> AliasScopes; |
| SmallVector<Metadata *> NoAliasList; |
| const size_t NumberVars = LDSVarsToTransform.size(); |
| if (NumberVars > 1) { |
| MDBuilder MDB(Ctx); |
| AliasScopes.reserve(NumberVars); |
| MDNode *Domain = MDB.createAnonymousAliasScopeDomain(); |
| for (size_t I = 0; I < NumberVars; I++) { |
| MDNode *Scope = MDB.createAnonymousAliasScope(Domain); |
| AliasScopes.push_back(Scope); |
| } |
| NoAliasList.append(&AliasScopes[1], AliasScopes.end()); |
| } |
| |
| // Replace uses of ith variable with a constantexpr to the corresponding |
| // field of the instance that will be allocated by AMDGPUMachineFunction |
| for (size_t I = 0; I < NumberVars; I++) { |
| GlobalVariable *GV = LDSVarsToTransform[I]; |
| Constant *GEP = Replacement.LDSVarsToConstantGEP[GV]; |
| |
| GV->replaceUsesWithIf(GEP, Predicate); |
| |
| APInt APOff(DL.getIndexTypeSizeInBits(GEP->getType()), 0); |
| GEP->stripAndAccumulateInBoundsConstantOffsets(DL, APOff); |
| uint64_t Offset = APOff.getZExtValue(); |
| |
| Align A = |
| commonAlignment(Replacement.SGV->getAlign().valueOrOne(), Offset); |
| |
| if (I) |
| NoAliasList[I - 1] = AliasScopes[I - 1]; |
| MDNode *NoAlias = |
| NoAliasList.empty() ? nullptr : MDNode::get(Ctx, NoAliasList); |
| MDNode *AliasScope = |
| AliasScopes.empty() ? nullptr : MDNode::get(Ctx, {AliasScopes[I]}); |
| |
| refineUsesAlignmentAndAA(GEP, A, DL, AliasScope, NoAlias); |
| } |
| } |
| |
| void refineUsesAlignmentAndAA(Value *Ptr, Align A, const DataLayout &DL, |
| MDNode *AliasScope, MDNode *NoAlias, |
| unsigned MaxDepth = 5) { |
| if (!MaxDepth || (A == 1 && !AliasScope)) |
| return; |
| |
| for (User *U : Ptr->users()) { |
| if (auto *I = dyn_cast<Instruction>(U)) { |
| if (AliasScope && I->mayReadOrWriteMemory()) { |
| MDNode *AS = I->getMetadata(LLVMContext::MD_alias_scope); |
| AS = (AS ? MDNode::getMostGenericAliasScope(AS, AliasScope) |
| : AliasScope); |
| I->setMetadata(LLVMContext::MD_alias_scope, AS); |
| |
| MDNode *NA = I->getMetadata(LLVMContext::MD_noalias); |
| NA = (NA ? MDNode::intersect(NA, NoAlias) : NoAlias); |
| I->setMetadata(LLVMContext::MD_noalias, NA); |
| } |
| } |
| |
| if (auto *LI = dyn_cast<LoadInst>(U)) { |
| LI->setAlignment(std::max(A, LI->getAlign())); |
| continue; |
| } |
| if (auto *SI = dyn_cast<StoreInst>(U)) { |
| if (SI->getPointerOperand() == Ptr) |
| SI->setAlignment(std::max(A, SI->getAlign())); |
| continue; |
| } |
| if (auto *AI = dyn_cast<AtomicRMWInst>(U)) { |
| // None of atomicrmw operations can work on pointers, but let's |
| // check it anyway in case it will or we will process ConstantExpr. |
| if (AI->getPointerOperand() == Ptr) |
| AI->setAlignment(std::max(A, AI->getAlign())); |
| continue; |
| } |
| if (auto *AI = dyn_cast<AtomicCmpXchgInst>(U)) { |
| if (AI->getPointerOperand() == Ptr) |
| AI->setAlignment(std::max(A, AI->getAlign())); |
| continue; |
| } |
| if (auto *GEP = dyn_cast<GetElementPtrInst>(U)) { |
| unsigned BitWidth = DL.getIndexTypeSizeInBits(GEP->getType()); |
| APInt Off(BitWidth, 0); |
| if (GEP->getPointerOperand() == Ptr) { |
| Align GA; |
| if (GEP->accumulateConstantOffset(DL, Off)) |
| GA = commonAlignment(A, Off.getLimitedValue()); |
| refineUsesAlignmentAndAA(GEP, GA, DL, AliasScope, NoAlias, |
| MaxDepth - 1); |
| } |
| continue; |
| } |
| if (auto *I = dyn_cast<Instruction>(U)) { |
| if (I->getOpcode() == Instruction::BitCast || |
| I->getOpcode() == Instruction::AddrSpaceCast) |
| refineUsesAlignmentAndAA(I, A, DL, AliasScope, NoAlias, MaxDepth - 1); |
| } |
| } |
| } |
| }; |
| |
| } // namespace |
| char AMDGPULowerModuleLDS::ID = 0; |
| |
| char &llvm::AMDGPULowerModuleLDSID = AMDGPULowerModuleLDS::ID; |
| |
| INITIALIZE_PASS(AMDGPULowerModuleLDS, DEBUG_TYPE, |
| "Lower uses of LDS variables from non-kernel functions", false, |
| false) |
| |
| ModulePass *llvm::createAMDGPULowerModuleLDSPass() { |
| return new AMDGPULowerModuleLDS(); |
| } |
| |
| PreservedAnalyses AMDGPULowerModuleLDSPass::run(Module &M, |
| ModuleAnalysisManager &) { |
| return AMDGPULowerModuleLDS().runOnModule(M) ? PreservedAnalyses::none() |
| : PreservedAnalyses::all(); |
| } |
