Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 1 | //===-- AMDGPUPromoteAlloca.cpp - Promote Allocas -------------------------===// |
| 2 | // |
Chandler Carruth | 2946cd7 | 2019-01-19 08:50:56 +0000 | [diff] [blame] | 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-exception |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 6 | // |
| 7 | //===----------------------------------------------------------------------===// |
| 8 | // |
| 9 | // This pass eliminates allocas by either converting them into vectors or |
| 10 | // by migrating them to local address space. |
| 11 | // |
| 12 | //===----------------------------------------------------------------------===// |
| 13 | |
| 14 | #include "AMDGPU.h" |
| 15 | #include "AMDGPUSubtarget.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 16 | #include "Utils/AMDGPUBaseInfo.h" |
| 17 | #include "llvm/ADT/APInt.h" |
| 18 | #include "llvm/ADT/None.h" |
| 19 | #include "llvm/ADT/STLExtras.h" |
| 20 | #include "llvm/ADT/StringRef.h" |
| 21 | #include "llvm/ADT/Triple.h" |
| 22 | #include "llvm/ADT/Twine.h" |
Changpeng Fang | c85abbd | 2017-01-24 19:06:28 +0000 | [diff] [blame] | 23 | #include "llvm/Analysis/CaptureTracking.h" |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 24 | #include "llvm/Analysis/ValueTracking.h" |
Francis Visoiu Mistrih | 8b61764 | 2017-05-18 17:21:13 +0000 | [diff] [blame] | 25 | #include "llvm/CodeGen/TargetPassConfig.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 26 | #include "llvm/IR/Attributes.h" |
| 27 | #include "llvm/IR/BasicBlock.h" |
| 28 | #include "llvm/IR/Constant.h" |
| 29 | #include "llvm/IR/Constants.h" |
| 30 | #include "llvm/IR/DataLayout.h" |
| 31 | #include "llvm/IR/DerivedTypes.h" |
| 32 | #include "llvm/IR/Function.h" |
| 33 | #include "llvm/IR/GlobalValue.h" |
| 34 | #include "llvm/IR/GlobalVariable.h" |
Chandler Carruth | 6bda14b | 2017-06-06 11:49:48 +0000 | [diff] [blame] | 35 | #include "llvm/IR/IRBuilder.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 36 | #include "llvm/IR/Instruction.h" |
| 37 | #include "llvm/IR/Instructions.h" |
Matt Arsenault | bafc9dc | 2016-03-11 08:20:50 +0000 | [diff] [blame] | 38 | #include "llvm/IR/IntrinsicInst.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 39 | #include "llvm/IR/Intrinsics.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 40 | #include "llvm/IR/LLVMContext.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 41 | #include "llvm/IR/Metadata.h" |
| 42 | #include "llvm/IR/Module.h" |
| 43 | #include "llvm/IR/Type.h" |
| 44 | #include "llvm/IR/User.h" |
| 45 | #include "llvm/IR/Value.h" |
| 46 | #include "llvm/Pass.h" |
| 47 | #include "llvm/Support/Casting.h" |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 48 | #include "llvm/Support/Debug.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 49 | #include "llvm/Support/ErrorHandling.h" |
| 50 | #include "llvm/Support/MathExtras.h" |
Benjamin Kramer | 16132e6 | 2015-03-23 18:07:13 +0000 | [diff] [blame] | 51 | #include "llvm/Support/raw_ostream.h" |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 52 | #include "llvm/Target/TargetMachine.h" |
| 53 | #include <algorithm> |
| 54 | #include <cassert> |
| 55 | #include <cstdint> |
| 56 | #include <map> |
| 57 | #include <tuple> |
| 58 | #include <utility> |
| 59 | #include <vector> |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 60 | |
| 61 | #define DEBUG_TYPE "amdgpu-promote-alloca" |
| 62 | |
| 63 | using namespace llvm; |
| 64 | |
| 65 | namespace { |
| 66 | |
Changpeng Fang | ba92059 | 2018-02-16 19:14:17 +0000 | [diff] [blame] | 67 | static cl::opt<bool> DisablePromoteAllocaToVector( |
| 68 | "disable-promote-alloca-to-vector", |
| 69 | cl::desc("Disable promote alloca to vector"), |
| 70 | cl::init(false)); |
| 71 | |
Yaxun Liu | 73bf0af | 2018-11-06 21:28:17 +0000 | [diff] [blame] | 72 | static cl::opt<bool> DisablePromoteAllocaToLDS( |
| 73 | "disable-promote-alloca-to-lds", |
| 74 | cl::desc("Disable promote alloca to LDS"), |
| 75 | cl::init(false)); |
| 76 | |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 77 | // FIXME: This can create globals so should be a module pass. |
Matt Arsenault | bafc9dc | 2016-03-11 08:20:50 +0000 | [diff] [blame] | 78 | class AMDGPUPromoteAlloca : public FunctionPass { |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 79 | private: |
| 80 | const TargetMachine *TM; |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 81 | Module *Mod = nullptr; |
| 82 | const DataLayout *DL = nullptr; |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 83 | |
| 84 | // FIXME: This should be per-kernel. |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 85 | uint32_t LocalMemLimit = 0; |
| 86 | uint32_t CurrentLocalMemUsage = 0; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 87 | |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 88 | bool IsAMDGCN = false; |
| 89 | bool IsAMDHSA = false; |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 90 | |
| 91 | std::pair<Value *, Value *> getLocalSizeYZ(IRBuilder<> &Builder); |
| 92 | Value *getWorkitemID(IRBuilder<> &Builder, unsigned N); |
| 93 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 94 | /// BaseAlloca is the alloca root the search started from. |
| 95 | /// Val may be that alloca or a recursive user of it. |
| 96 | bool collectUsesWithPtrTypes(Value *BaseAlloca, |
| 97 | Value *Val, |
| 98 | std::vector<Value*> &WorkList) const; |
| 99 | |
| 100 | /// Val is a derived pointer from Alloca. OpIdx0/OpIdx1 are the operand |
| 101 | /// indices to an instruction with 2 pointer inputs (e.g. select, icmp). |
| 102 | /// Returns true if both operands are derived from the same alloca. Val should |
| 103 | /// be the same value as one of the input operands of UseInst. |
| 104 | bool binaryOpIsDerivedFromSameAlloca(Value *Alloca, Value *Val, |
| 105 | Instruction *UseInst, |
| 106 | int OpIdx0, int OpIdx1) const; |
| 107 | |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 108 | /// Check whether we have enough local memory for promotion. |
| 109 | bool hasSufficientLocalMem(const Function &F); |
| 110 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 111 | public: |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 112 | static char ID; |
| 113 | |
Francis Visoiu Mistrih | 8b61764 | 2017-05-18 17:21:13 +0000 | [diff] [blame] | 114 | AMDGPUPromoteAlloca() : FunctionPass(ID) {} |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 115 | |
Benjamin Kramer | 8c90fd7 | 2014-09-03 11:41:21 +0000 | [diff] [blame] | 116 | bool doInitialization(Module &M) override; |
| 117 | bool runOnFunction(Function &F) override; |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 118 | |
Mehdi Amini | 117296c | 2016-10-01 02:56:57 +0000 | [diff] [blame] | 119 | StringRef getPassName() const override { return "AMDGPU Promote Alloca"; } |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 120 | |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 121 | bool handleAlloca(AllocaInst &I, bool SufficientLDS); |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 122 | |
| 123 | void getAnalysisUsage(AnalysisUsage &AU) const override { |
| 124 | AU.setPreservesCFG(); |
| 125 | FunctionPass::getAnalysisUsage(AU); |
| 126 | } |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 127 | }; |
| 128 | |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 129 | } // end anonymous namespace |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 130 | |
| 131 | char AMDGPUPromoteAlloca::ID = 0; |
| 132 | |
Francis Visoiu Mistrih | 8b61764 | 2017-05-18 17:21:13 +0000 | [diff] [blame] | 133 | INITIALIZE_PASS(AMDGPUPromoteAlloca, DEBUG_TYPE, |
| 134 | "AMDGPU promote alloca to vector or LDS", false, false) |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 135 | |
| 136 | char &llvm::AMDGPUPromoteAllocaID = AMDGPUPromoteAlloca::ID; |
| 137 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 138 | bool AMDGPUPromoteAlloca::doInitialization(Module &M) { |
| 139 | Mod = &M; |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 140 | DL = &Mod->getDataLayout(); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 141 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 142 | return false; |
| 143 | } |
| 144 | |
| 145 | bool AMDGPUPromoteAlloca::runOnFunction(Function &F) { |
Francis Visoiu Mistrih | 8b61764 | 2017-05-18 17:21:13 +0000 | [diff] [blame] | 146 | if (skipFunction(F)) |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 147 | return false; |
| 148 | |
Francis Visoiu Mistrih | 8b61764 | 2017-05-18 17:21:13 +0000 | [diff] [blame] | 149 | if (auto *TPC = getAnalysisIfAvailable<TargetPassConfig>()) |
| 150 | TM = &TPC->getTM<TargetMachine>(); |
| 151 | else |
| 152 | return false; |
| 153 | |
| 154 | const Triple &TT = TM->getTargetTriple(); |
| 155 | IsAMDGCN = TT.getArch() == Triple::amdgcn; |
| 156 | IsAMDHSA = TT.getOS() == Triple::AMDHSA; |
| 157 | |
Tom Stellard | 5bfbae5 | 2018-07-11 20:59:01 +0000 | [diff] [blame] | 158 | const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(*TM, F); |
Matt Arsenault | 03d8584 | 2016-06-27 20:32:13 +0000 | [diff] [blame] | 159 | if (!ST.isPromoteAllocaEnabled()) |
| 160 | return false; |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 161 | |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 162 | bool SufficientLDS = hasSufficientLocalMem(F); |
| 163 | bool Changed = false; |
Matt Arsenault | bafc9dc | 2016-03-11 08:20:50 +0000 | [diff] [blame] | 164 | BasicBlock &EntryBB = *F.begin(); |
| 165 | for (auto I = EntryBB.begin(), E = EntryBB.end(); I != E; ) { |
| 166 | AllocaInst *AI = dyn_cast<AllocaInst>(I); |
| 167 | |
| 168 | ++I; |
| 169 | if (AI) |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 170 | Changed |= handleAlloca(*AI, SufficientLDS); |
Matt Arsenault | bafc9dc | 2016-03-11 08:20:50 +0000 | [diff] [blame] | 171 | } |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 172 | |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 173 | return Changed; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 174 | } |
| 175 | |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 176 | std::pair<Value *, Value *> |
| 177 | AMDGPUPromoteAlloca::getLocalSizeYZ(IRBuilder<> &Builder) { |
Tom Stellard | c5a154d | 2018-06-28 23:47:12 +0000 | [diff] [blame] | 178 | const Function &F = *Builder.GetInsertBlock()->getParent(); |
Tom Stellard | 5bfbae5 | 2018-07-11 20:59:01 +0000 | [diff] [blame] | 179 | const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(*TM, F); |
Stanislav Mekhanoshin | c90347d | 2017-04-12 20:48:56 +0000 | [diff] [blame] | 180 | |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 181 | if (!IsAMDHSA) { |
| 182 | Function *LocalSizeYFn |
| 183 | = Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_y); |
| 184 | Function *LocalSizeZFn |
| 185 | = Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_z); |
| 186 | |
| 187 | CallInst *LocalSizeY = Builder.CreateCall(LocalSizeYFn, {}); |
| 188 | CallInst *LocalSizeZ = Builder.CreateCall(LocalSizeZFn, {}); |
| 189 | |
Stanislav Mekhanoshin | c90347d | 2017-04-12 20:48:56 +0000 | [diff] [blame] | 190 | ST.makeLIDRangeMetadata(LocalSizeY); |
| 191 | ST.makeLIDRangeMetadata(LocalSizeZ); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 192 | |
| 193 | return std::make_pair(LocalSizeY, LocalSizeZ); |
| 194 | } |
| 195 | |
| 196 | // We must read the size out of the dispatch pointer. |
| 197 | assert(IsAMDGCN); |
| 198 | |
| 199 | // We are indexing into this struct, and want to extract the workgroup_size_* |
| 200 | // fields. |
| 201 | // |
| 202 | // typedef struct hsa_kernel_dispatch_packet_s { |
| 203 | // uint16_t header; |
| 204 | // uint16_t setup; |
| 205 | // uint16_t workgroup_size_x ; |
| 206 | // uint16_t workgroup_size_y; |
| 207 | // uint16_t workgroup_size_z; |
| 208 | // uint16_t reserved0; |
| 209 | // uint32_t grid_size_x ; |
| 210 | // uint32_t grid_size_y ; |
| 211 | // uint32_t grid_size_z; |
| 212 | // |
| 213 | // uint32_t private_segment_size; |
| 214 | // uint32_t group_segment_size; |
| 215 | // uint64_t kernel_object; |
| 216 | // |
| 217 | // #ifdef HSA_LARGE_MODEL |
| 218 | // void *kernarg_address; |
| 219 | // #elif defined HSA_LITTLE_ENDIAN |
| 220 | // void *kernarg_address; |
| 221 | // uint32_t reserved1; |
| 222 | // #else |
| 223 | // uint32_t reserved1; |
| 224 | // void *kernarg_address; |
| 225 | // #endif |
| 226 | // uint64_t reserved2; |
| 227 | // hsa_signal_t completion_signal; // uint64_t wrapper |
| 228 | // } hsa_kernel_dispatch_packet_t |
| 229 | // |
| 230 | Function *DispatchPtrFn |
| 231 | = Intrinsic::getDeclaration(Mod, Intrinsic::amdgcn_dispatch_ptr); |
| 232 | |
| 233 | CallInst *DispatchPtr = Builder.CreateCall(DispatchPtrFn, {}); |
Reid Kleckner | b518054 | 2017-03-21 16:57:19 +0000 | [diff] [blame] | 234 | DispatchPtr->addAttribute(AttributeList::ReturnIndex, Attribute::NoAlias); |
| 235 | DispatchPtr->addAttribute(AttributeList::ReturnIndex, Attribute::NonNull); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 236 | |
| 237 | // Size of the dispatch packet struct. |
Reid Kleckner | b518054 | 2017-03-21 16:57:19 +0000 | [diff] [blame] | 238 | DispatchPtr->addDereferenceableAttr(AttributeList::ReturnIndex, 64); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 239 | |
| 240 | Type *I32Ty = Type::getInt32Ty(Mod->getContext()); |
| 241 | Value *CastDispatchPtr = Builder.CreateBitCast( |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 242 | DispatchPtr, PointerType::get(I32Ty, AMDGPUAS::CONSTANT_ADDRESS)); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 243 | |
| 244 | // We could do a single 64-bit load here, but it's likely that the basic |
| 245 | // 32-bit and extract sequence is already present, and it is probably easier |
| 246 | // to CSE this. The loads should be mergable later anyway. |
James Y Knight | 7716075 | 2019-02-01 20:44:47 +0000 | [diff] [blame] | 247 | Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(I32Ty, CastDispatchPtr, 1); |
James Y Knight | 14359ef | 2019-02-01 20:44:24 +0000 | [diff] [blame] | 248 | LoadInst *LoadXY = Builder.CreateAlignedLoad(I32Ty, GEPXY, 4); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 249 | |
James Y Knight | 7716075 | 2019-02-01 20:44:47 +0000 | [diff] [blame] | 250 | Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(I32Ty, CastDispatchPtr, 2); |
James Y Knight | 14359ef | 2019-02-01 20:44:24 +0000 | [diff] [blame] | 251 | LoadInst *LoadZU = Builder.CreateAlignedLoad(I32Ty, GEPZU, 4); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 252 | |
Eugene Zelenko | 734bb7b | 2017-01-20 17:52:16 +0000 | [diff] [blame] | 253 | MDNode *MD = MDNode::get(Mod->getContext(), None); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 254 | LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD); |
| 255 | LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD); |
Stanislav Mekhanoshin | c90347d | 2017-04-12 20:48:56 +0000 | [diff] [blame] | 256 | ST.makeLIDRangeMetadata(LoadZU); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 257 | |
| 258 | // Extract y component. Upper half of LoadZU should be zero already. |
| 259 | Value *Y = Builder.CreateLShr(LoadXY, 16); |
| 260 | |
| 261 | return std::make_pair(Y, LoadZU); |
| 262 | } |
| 263 | |
| 264 | Value *AMDGPUPromoteAlloca::getWorkitemID(IRBuilder<> &Builder, unsigned N) { |
Tom Stellard | 5bfbae5 | 2018-07-11 20:59:01 +0000 | [diff] [blame] | 265 | const AMDGPUSubtarget &ST = |
| 266 | AMDGPUSubtarget::get(*TM, *Builder.GetInsertBlock()->getParent()); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 267 | Intrinsic::ID IntrID = Intrinsic::ID::not_intrinsic; |
| 268 | |
| 269 | switch (N) { |
| 270 | case 0: |
| 271 | IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_x |
| 272 | : Intrinsic::r600_read_tidig_x; |
| 273 | break; |
| 274 | case 1: |
| 275 | IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_y |
| 276 | : Intrinsic::r600_read_tidig_y; |
| 277 | break; |
| 278 | |
| 279 | case 2: |
| 280 | IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_z |
| 281 | : Intrinsic::r600_read_tidig_z; |
| 282 | break; |
| 283 | default: |
| 284 | llvm_unreachable("invalid dimension"); |
| 285 | } |
| 286 | |
| 287 | Function *WorkitemIdFn = Intrinsic::getDeclaration(Mod, IntrID); |
| 288 | CallInst *CI = Builder.CreateCall(WorkitemIdFn); |
Stanislav Mekhanoshin | c90347d | 2017-04-12 20:48:56 +0000 | [diff] [blame] | 289 | ST.makeLIDRangeMetadata(CI); |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 290 | |
| 291 | return CI; |
| 292 | } |
| 293 | |
Matt Arsenault | 37ab4cf | 2017-09-14 18:02:29 +0000 | [diff] [blame] | 294 | static VectorType *arrayTypeToVecType(ArrayType *ArrayTy) { |
| 295 | return VectorType::get(ArrayTy->getElementType(), |
| 296 | ArrayTy->getNumElements()); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 297 | } |
| 298 | |
Benjamin Kramer | c6cc58e | 2014-10-04 16:55:56 +0000 | [diff] [blame] | 299 | static Value * |
| 300 | calculateVectorIndex(Value *Ptr, |
| 301 | const std::map<GetElementPtrInst *, Value *> &GEPIdx) { |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 302 | GetElementPtrInst *GEP = cast<GetElementPtrInst>(Ptr); |
| 303 | |
Benjamin Kramer | c6cc58e | 2014-10-04 16:55:56 +0000 | [diff] [blame] | 304 | auto I = GEPIdx.find(GEP); |
| 305 | return I == GEPIdx.end() ? nullptr : I->second; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 306 | } |
| 307 | |
| 308 | static Value* GEPToVectorIndex(GetElementPtrInst *GEP) { |
| 309 | // FIXME we only support simple cases |
| 310 | if (GEP->getNumOperands() != 3) |
Matt Arsenault | efb2454 | 2016-07-18 18:34:53 +0000 | [diff] [blame] | 311 | return nullptr; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 312 | |
| 313 | ConstantInt *I0 = dyn_cast<ConstantInt>(GEP->getOperand(1)); |
| 314 | if (!I0 || !I0->isZero()) |
Matt Arsenault | efb2454 | 2016-07-18 18:34:53 +0000 | [diff] [blame] | 315 | return nullptr; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 316 | |
| 317 | return GEP->getOperand(2); |
| 318 | } |
| 319 | |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 320 | // Not an instruction handled below to turn into a vector. |
| 321 | // |
| 322 | // TODO: Check isTriviallyVectorizable for calls and handle other |
| 323 | // instructions. |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 324 | static bool canVectorizeInst(Instruction *Inst, User *User) { |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 325 | switch (Inst->getOpcode()) { |
Changpeng Fang | 161e8c39af | 2017-05-12 20:31:12 +0000 | [diff] [blame] | 326 | case Instruction::Load: { |
Changpeng Fang | 860d460 | 2018-05-17 21:49:44 +0000 | [diff] [blame] | 327 | // Currently only handle the case where the Pointer Operand is a GEP. |
| 328 | // Also we could not vectorize volatile or atomic loads. |
Changpeng Fang | 161e8c39af | 2017-05-12 20:31:12 +0000 | [diff] [blame] | 329 | LoadInst *LI = cast<LoadInst>(Inst); |
Stanislav Mekhanoshin | 6cc8b2f | 2018-11-08 00:16:23 +0000 | [diff] [blame] | 330 | if (isa<AllocaInst>(User) && |
| 331 | LI->getPointerOperandType() == User->getType() && |
| 332 | isa<VectorType>(LI->getType())) |
| 333 | return true; |
Changpeng Fang | 860d460 | 2018-05-17 21:49:44 +0000 | [diff] [blame] | 334 | return isa<GetElementPtrInst>(LI->getPointerOperand()) && LI->isSimple(); |
Changpeng Fang | 161e8c39af | 2017-05-12 20:31:12 +0000 | [diff] [blame] | 335 | } |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 336 | case Instruction::BitCast: |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 337 | return true; |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 338 | case Instruction::Store: { |
David Stuttard | 82618ba | 2017-06-09 14:16:22 +0000 | [diff] [blame] | 339 | // Must be the stored pointer operand, not a stored value, plus |
| 340 | // since it should be canonical form, the User should be a GEP. |
Changpeng Fang | 860d460 | 2018-05-17 21:49:44 +0000 | [diff] [blame] | 341 | // Also we could not vectorize volatile or atomic stores. |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 342 | StoreInst *SI = cast<StoreInst>(Inst); |
Stanislav Mekhanoshin | 6cc8b2f | 2018-11-08 00:16:23 +0000 | [diff] [blame] | 343 | if (isa<AllocaInst>(User) && |
| 344 | SI->getPointerOperandType() == User->getType() && |
| 345 | isa<VectorType>(SI->getValueOperand()->getType())) |
| 346 | return true; |
Changpeng Fang | 860d460 | 2018-05-17 21:49:44 +0000 | [diff] [blame] | 347 | return (SI->getPointerOperand() == User) && isa<GetElementPtrInst>(User) && SI->isSimple(); |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 348 | } |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 349 | default: |
| 350 | return false; |
| 351 | } |
| 352 | } |
| 353 | |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 354 | static bool tryPromoteAllocaToVector(AllocaInst *Alloca) { |
Changpeng Fang | ba92059 | 2018-02-16 19:14:17 +0000 | [diff] [blame] | 355 | |
| 356 | if (DisablePromoteAllocaToVector) { |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 357 | LLVM_DEBUG(dbgs() << " Promotion alloca to vector is disabled\n"); |
Changpeng Fang | ba92059 | 2018-02-16 19:14:17 +0000 | [diff] [blame] | 358 | return false; |
| 359 | } |
| 360 | |
Stanislav Mekhanoshin | 6cc8b2f | 2018-11-08 00:16:23 +0000 | [diff] [blame] | 361 | Type *AT = Alloca->getAllocatedType(); |
| 362 | SequentialType *AllocaTy = dyn_cast<SequentialType>(AT); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 363 | |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 364 | LLVM_DEBUG(dbgs() << "Alloca candidate for vectorization\n"); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 365 | |
| 366 | // FIXME: There is no reason why we can't support larger arrays, we |
| 367 | // are just being conservative for now. |
David Stuttard | 82618ba | 2017-06-09 14:16:22 +0000 | [diff] [blame] | 368 | // FIXME: We also reject alloca's of the form [ 2 x [ 2 x i32 ]] or equivalent. Potentially these |
| 369 | // could also be promoted but we don't currently handle this case |
Matt Arsenault | fb8cdba | 2016-02-02 19:32:35 +0000 | [diff] [blame] | 370 | if (!AllocaTy || |
Changpeng Fang | ba92059 | 2018-02-16 19:14:17 +0000 | [diff] [blame] | 371 | AllocaTy->getNumElements() > 16 || |
Matt Arsenault | 37ab4cf | 2017-09-14 18:02:29 +0000 | [diff] [blame] | 372 | AllocaTy->getNumElements() < 2 || |
| 373 | !VectorType::isValidElementType(AllocaTy->getElementType())) { |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 374 | LLVM_DEBUG(dbgs() << " Cannot convert type to vector\n"); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 375 | return false; |
| 376 | } |
| 377 | |
| 378 | std::map<GetElementPtrInst*, Value*> GEPVectorIdx; |
| 379 | std::vector<Value*> WorkList; |
| 380 | for (User *AllocaUser : Alloca->users()) { |
| 381 | GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(AllocaUser); |
| 382 | if (!GEP) { |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 383 | if (!canVectorizeInst(cast<Instruction>(AllocaUser), Alloca)) |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 384 | return false; |
| 385 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 386 | WorkList.push_back(AllocaUser); |
| 387 | continue; |
| 388 | } |
| 389 | |
| 390 | Value *Index = GEPToVectorIndex(GEP); |
| 391 | |
| 392 | // If we can't compute a vector index from this GEP, then we can't |
| 393 | // promote this alloca to vector. |
| 394 | if (!Index) { |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 395 | LLVM_DEBUG(dbgs() << " Cannot compute vector index for GEP " << *GEP |
| 396 | << '\n'); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 397 | return false; |
| 398 | } |
| 399 | |
| 400 | GEPVectorIdx[GEP] = Index; |
| 401 | for (User *GEPUser : AllocaUser->users()) { |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 402 | if (!canVectorizeInst(cast<Instruction>(GEPUser), AllocaUser)) |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 403 | return false; |
| 404 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 405 | WorkList.push_back(GEPUser); |
| 406 | } |
| 407 | } |
| 408 | |
Stanislav Mekhanoshin | 6cc8b2f | 2018-11-08 00:16:23 +0000 | [diff] [blame] | 409 | VectorType *VectorTy = dyn_cast<VectorType>(AllocaTy); |
| 410 | if (!VectorTy) |
| 411 | VectorTy = arrayTypeToVecType(cast<ArrayType>(AllocaTy)); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 412 | |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 413 | LLVM_DEBUG(dbgs() << " Converting alloca to vector " << *AllocaTy << " -> " |
| 414 | << *VectorTy << '\n'); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 415 | |
Matt Arsenault | fb8cdba | 2016-02-02 19:32:35 +0000 | [diff] [blame] | 416 | for (Value *V : WorkList) { |
| 417 | Instruction *Inst = cast<Instruction>(V); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 418 | IRBuilder<> Builder(Inst); |
| 419 | switch (Inst->getOpcode()) { |
| 420 | case Instruction::Load: { |
Stanislav Mekhanoshin | 6cc8b2f | 2018-11-08 00:16:23 +0000 | [diff] [blame] | 421 | if (Inst->getType() == AT) |
| 422 | break; |
| 423 | |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 424 | Type *VecPtrTy = VectorTy->getPointerTo(AMDGPUAS::PRIVATE_ADDRESS); |
David Stuttard | 82618ba | 2017-06-09 14:16:22 +0000 | [diff] [blame] | 425 | Value *Ptr = cast<LoadInst>(Inst)->getPointerOperand(); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 426 | Value *Index = calculateVectorIndex(Ptr, GEPVectorIdx); |
Matt Arsenault | efb2454 | 2016-07-18 18:34:53 +0000 | [diff] [blame] | 427 | |
| 428 | Value *BitCast = Builder.CreateBitCast(Alloca, VecPtrTy); |
James Y Knight | 14359ef | 2019-02-01 20:44:24 +0000 | [diff] [blame] | 429 | Value *VecValue = Builder.CreateLoad(VectorTy, BitCast); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 430 | Value *ExtractElement = Builder.CreateExtractElement(VecValue, Index); |
| 431 | Inst->replaceAllUsesWith(ExtractElement); |
| 432 | Inst->eraseFromParent(); |
| 433 | break; |
| 434 | } |
| 435 | case Instruction::Store: { |
David Stuttard | 82618ba | 2017-06-09 14:16:22 +0000 | [diff] [blame] | 436 | StoreInst *SI = cast<StoreInst>(Inst); |
Stanislav Mekhanoshin | 6cc8b2f | 2018-11-08 00:16:23 +0000 | [diff] [blame] | 437 | if (SI->getValueOperand()->getType() == AT) |
| 438 | break; |
| 439 | |
| 440 | Type *VecPtrTy = VectorTy->getPointerTo(AMDGPUAS::PRIVATE_ADDRESS); |
David Stuttard | 82618ba | 2017-06-09 14:16:22 +0000 | [diff] [blame] | 441 | Value *Ptr = SI->getPointerOperand(); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 442 | Value *Index = calculateVectorIndex(Ptr, GEPVectorIdx); |
Matt Arsenault | efb2454 | 2016-07-18 18:34:53 +0000 | [diff] [blame] | 443 | Value *BitCast = Builder.CreateBitCast(Alloca, VecPtrTy); |
James Y Knight | 14359ef | 2019-02-01 20:44:24 +0000 | [diff] [blame] | 444 | Value *VecValue = Builder.CreateLoad(VectorTy, BitCast); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 445 | Value *NewVecValue = Builder.CreateInsertElement(VecValue, |
David Stuttard | 82618ba | 2017-06-09 14:16:22 +0000 | [diff] [blame] | 446 | SI->getValueOperand(), |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 447 | Index); |
| 448 | Builder.CreateStore(NewVecValue, BitCast); |
| 449 | Inst->eraseFromParent(); |
| 450 | break; |
| 451 | } |
| 452 | case Instruction::BitCast: |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 453 | case Instruction::AddrSpaceCast: |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 454 | break; |
| 455 | |
| 456 | default: |
Matt Arsenault | 642d2e7 | 2014-06-27 16:52:49 +0000 | [diff] [blame] | 457 | llvm_unreachable("Inconsistency in instructions promotable to vector"); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 458 | } |
| 459 | } |
| 460 | return true; |
| 461 | } |
| 462 | |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 463 | static bool isCallPromotable(CallInst *CI) { |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 464 | IntrinsicInst *II = dyn_cast<IntrinsicInst>(CI); |
| 465 | if (!II) |
| 466 | return false; |
| 467 | |
| 468 | switch (II->getIntrinsicID()) { |
| 469 | case Intrinsic::memcpy: |
Matt Arsenault | 7e747f1 | 2016-02-02 20:28:10 +0000 | [diff] [blame] | 470 | case Intrinsic::memmove: |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 471 | case Intrinsic::memset: |
| 472 | case Intrinsic::lifetime_start: |
| 473 | case Intrinsic::lifetime_end: |
| 474 | case Intrinsic::invariant_start: |
| 475 | case Intrinsic::invariant_end: |
Piotr Padlewski | 5dde809 | 2018-05-03 11:03:01 +0000 | [diff] [blame] | 476 | case Intrinsic::launder_invariant_group: |
Piotr Padlewski | 5b3db45 | 2018-07-02 04:49:30 +0000 | [diff] [blame] | 477 | case Intrinsic::strip_invariant_group: |
Matt Arsenault | 7e747f1 | 2016-02-02 20:28:10 +0000 | [diff] [blame] | 478 | case Intrinsic::objectsize: |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 479 | return true; |
| 480 | default: |
| 481 | return false; |
| 482 | } |
| 483 | } |
| 484 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 485 | bool AMDGPUPromoteAlloca::binaryOpIsDerivedFromSameAlloca(Value *BaseAlloca, |
| 486 | Value *Val, |
| 487 | Instruction *Inst, |
| 488 | int OpIdx0, |
| 489 | int OpIdx1) const { |
| 490 | // Figure out which operand is the one we might not be promoting. |
| 491 | Value *OtherOp = Inst->getOperand(OpIdx0); |
| 492 | if (Val == OtherOp) |
| 493 | OtherOp = Inst->getOperand(OpIdx1); |
| 494 | |
Matt Arsenault | 891fccc | 2016-05-18 15:57:21 +0000 | [diff] [blame] | 495 | if (isa<ConstantPointerNull>(OtherOp)) |
| 496 | return true; |
| 497 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 498 | Value *OtherObj = GetUnderlyingObject(OtherOp, *DL); |
| 499 | if (!isa<AllocaInst>(OtherObj)) |
| 500 | return false; |
| 501 | |
| 502 | // TODO: We should be able to replace undefs with the right pointer type. |
| 503 | |
| 504 | // TODO: If we know the other base object is another promotable |
| 505 | // alloca, not necessarily this alloca, we can do this. The |
| 506 | // important part is both must have the same address space at |
| 507 | // the end. |
| 508 | if (OtherObj != BaseAlloca) { |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 509 | LLVM_DEBUG( |
| 510 | dbgs() << "Found a binary instruction with another alloca object\n"); |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 511 | return false; |
| 512 | } |
| 513 | |
| 514 | return true; |
| 515 | } |
| 516 | |
| 517 | bool AMDGPUPromoteAlloca::collectUsesWithPtrTypes( |
| 518 | Value *BaseAlloca, |
| 519 | Value *Val, |
| 520 | std::vector<Value*> &WorkList) const { |
| 521 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 522 | for (User *User : Val->users()) { |
David Majnemer | 0d955d0 | 2016-08-11 22:21:41 +0000 | [diff] [blame] | 523 | if (is_contained(WorkList, User)) |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 524 | continue; |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 525 | |
Matt Arsenault | fdcd39a | 2015-07-28 18:29:14 +0000 | [diff] [blame] | 526 | if (CallInst *CI = dyn_cast<CallInst>(User)) { |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 527 | if (!isCallPromotable(CI)) |
Matt Arsenault | fdcd39a | 2015-07-28 18:29:14 +0000 | [diff] [blame] | 528 | return false; |
| 529 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 530 | WorkList.push_back(User); |
| 531 | continue; |
| 532 | } |
Tom Stellard | 5b2927f | 2014-10-31 20:52:04 +0000 | [diff] [blame] | 533 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 534 | Instruction *UseInst = cast<Instruction>(User); |
| 535 | if (UseInst->getOpcode() == Instruction::PtrToInt) |
Tom Stellard | 5b2927f | 2014-10-31 20:52:04 +0000 | [diff] [blame] | 536 | return false; |
| 537 | |
Matt Arsenault | 210b7cf | 2016-07-18 19:00:07 +0000 | [diff] [blame] | 538 | if (LoadInst *LI = dyn_cast<LoadInst>(UseInst)) { |
Matt Arsenault | c438ef5 | 2016-05-18 23:20:24 +0000 | [diff] [blame] | 539 | if (LI->isVolatile()) |
| 540 | return false; |
| 541 | |
| 542 | continue; |
| 543 | } |
| 544 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 545 | if (StoreInst *SI = dyn_cast<StoreInst>(UseInst)) { |
Matt Arsenault | 0a30e45 | 2016-03-23 23:17:29 +0000 | [diff] [blame] | 546 | if (SI->isVolatile()) |
| 547 | return false; |
| 548 | |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 549 | // Reject if the stored value is not the pointer operand. |
| 550 | if (SI->getPointerOperand() != Val) |
| 551 | return false; |
Matt Arsenault | 210b7cf | 2016-07-18 19:00:07 +0000 | [diff] [blame] | 552 | } else if (AtomicRMWInst *RMW = dyn_cast<AtomicRMWInst>(UseInst)) { |
Matt Arsenault | 0a30e45 | 2016-03-23 23:17:29 +0000 | [diff] [blame] | 553 | if (RMW->isVolatile()) |
| 554 | return false; |
Matt Arsenault | 210b7cf | 2016-07-18 19:00:07 +0000 | [diff] [blame] | 555 | } else if (AtomicCmpXchgInst *CAS = dyn_cast<AtomicCmpXchgInst>(UseInst)) { |
Matt Arsenault | 0a30e45 | 2016-03-23 23:17:29 +0000 | [diff] [blame] | 556 | if (CAS->isVolatile()) |
| 557 | return false; |
Matt Arsenault | 7227cc1 | 2015-07-28 18:47:00 +0000 | [diff] [blame] | 558 | } |
| 559 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 560 | // Only promote a select if we know that the other select operand |
| 561 | // is from another pointer that will also be promoted. |
| 562 | if (ICmpInst *ICmp = dyn_cast<ICmpInst>(UseInst)) { |
| 563 | if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, ICmp, 0, 1)) |
| 564 | return false; |
Matt Arsenault | 891fccc | 2016-05-18 15:57:21 +0000 | [diff] [blame] | 565 | |
| 566 | // May need to rewrite constant operands. |
| 567 | WorkList.push_back(ICmp); |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 568 | } |
| 569 | |
Matt Arsenault | 2402b95 | 2016-12-10 00:52:50 +0000 | [diff] [blame] | 570 | if (UseInst->getOpcode() == Instruction::AddrSpaceCast) { |
Changpeng Fang | c85abbd | 2017-01-24 19:06:28 +0000 | [diff] [blame] | 571 | // Give up if the pointer may be captured. |
| 572 | if (PointerMayBeCaptured(UseInst, true, true)) |
| 573 | return false; |
Matt Arsenault | 2402b95 | 2016-12-10 00:52:50 +0000 | [diff] [blame] | 574 | // Don't collect the users of this. |
| 575 | WorkList.push_back(User); |
| 576 | continue; |
| 577 | } |
| 578 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 579 | if (!User->getType()->isPointerTy()) |
| 580 | continue; |
Tom Stellard | 5b2927f | 2014-10-31 20:52:04 +0000 | [diff] [blame] | 581 | |
Matt Arsenault | de42081 | 2016-02-02 21:16:12 +0000 | [diff] [blame] | 582 | if (GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(UseInst)) { |
| 583 | // Be conservative if an address could be computed outside the bounds of |
| 584 | // the alloca. |
| 585 | if (!GEP->isInBounds()) |
| 586 | return false; |
| 587 | } |
| 588 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 589 | // Only promote a select if we know that the other select operand is from |
| 590 | // another pointer that will also be promoted. |
| 591 | if (SelectInst *SI = dyn_cast<SelectInst>(UseInst)) { |
| 592 | if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, SI, 1, 2)) |
| 593 | return false; |
| 594 | } |
| 595 | |
| 596 | // Repeat for phis. |
| 597 | if (PHINode *Phi = dyn_cast<PHINode>(UseInst)) { |
| 598 | // TODO: Handle more complex cases. We should be able to replace loops |
| 599 | // over arrays. |
| 600 | switch (Phi->getNumIncomingValues()) { |
| 601 | case 1: |
| 602 | break; |
| 603 | case 2: |
| 604 | if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, Phi, 0, 1)) |
| 605 | return false; |
| 606 | break; |
| 607 | default: |
| 608 | return false; |
| 609 | } |
| 610 | } |
| 611 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 612 | WorkList.push_back(User); |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 613 | if (!collectUsesWithPtrTypes(BaseAlloca, User, WorkList)) |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 614 | return false; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 615 | } |
Matt Arsenault | ad13484 | 2016-02-02 19:18:53 +0000 | [diff] [blame] | 616 | |
| 617 | return true; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 618 | } |
| 619 | |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 620 | bool AMDGPUPromoteAlloca::hasSufficientLocalMem(const Function &F) { |
| 621 | |
| 622 | FunctionType *FTy = F.getFunctionType(); |
Tom Stellard | 5bfbae5 | 2018-07-11 20:59:01 +0000 | [diff] [blame] | 623 | const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(*TM, F); |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 624 | |
| 625 | // If the function has any arguments in the local address space, then it's |
| 626 | // possible these arguments require the entire local memory space, so |
| 627 | // we cannot use local memory in the pass. |
| 628 | for (Type *ParamTy : FTy->params()) { |
| 629 | PointerType *PtrTy = dyn_cast<PointerType>(ParamTy); |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 630 | if (PtrTy && PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) { |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 631 | LocalMemLimit = 0; |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 632 | LLVM_DEBUG(dbgs() << "Function has local memory argument. Promoting to " |
| 633 | "local memory disabled.\n"); |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 634 | return false; |
| 635 | } |
| 636 | } |
| 637 | |
| 638 | LocalMemLimit = ST.getLocalMemorySize(); |
| 639 | if (LocalMemLimit == 0) |
| 640 | return false; |
| 641 | |
| 642 | const DataLayout &DL = Mod->getDataLayout(); |
| 643 | |
| 644 | // Check how much local memory is being used by global objects |
| 645 | CurrentLocalMemUsage = 0; |
| 646 | for (GlobalVariable &GV : Mod->globals()) { |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 647 | if (GV.getType()->getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 648 | continue; |
| 649 | |
| 650 | for (const User *U : GV.users()) { |
| 651 | const Instruction *Use = dyn_cast<Instruction>(U); |
| 652 | if (!Use) |
| 653 | continue; |
| 654 | |
| 655 | if (Use->getParent()->getParent() == &F) { |
| 656 | unsigned Align = GV.getAlignment(); |
| 657 | if (Align == 0) |
| 658 | Align = DL.getABITypeAlignment(GV.getValueType()); |
| 659 | |
| 660 | // FIXME: Try to account for padding here. The padding is currently |
| 661 | // determined from the inverse order of uses in the function. I'm not |
| 662 | // sure if the use list order is in any way connected to this, so the |
| 663 | // total reported size is likely incorrect. |
| 664 | uint64_t AllocSize = DL.getTypeAllocSize(GV.getValueType()); |
| 665 | CurrentLocalMemUsage = alignTo(CurrentLocalMemUsage, Align); |
| 666 | CurrentLocalMemUsage += AllocSize; |
| 667 | break; |
| 668 | } |
| 669 | } |
| 670 | } |
| 671 | |
| 672 | unsigned MaxOccupancy = ST.getOccupancyWithLocalMemSize(CurrentLocalMemUsage, |
| 673 | F); |
| 674 | |
| 675 | // Restrict local memory usage so that we don't drastically reduce occupancy, |
| 676 | // unless it is already significantly reduced. |
| 677 | |
| 678 | // TODO: Have some sort of hint or other heuristics to guess occupancy based |
| 679 | // on other factors.. |
| 680 | unsigned OccupancyHint = ST.getWavesPerEU(F).second; |
| 681 | if (OccupancyHint == 0) |
| 682 | OccupancyHint = 7; |
| 683 | |
| 684 | // Clamp to max value. |
| 685 | OccupancyHint = std::min(OccupancyHint, ST.getMaxWavesPerEU()); |
| 686 | |
| 687 | // Check the hint but ignore it if it's obviously wrong from the existing LDS |
| 688 | // usage. |
| 689 | MaxOccupancy = std::min(OccupancyHint, MaxOccupancy); |
| 690 | |
| 691 | |
| 692 | // Round up to the next tier of usage. |
| 693 | unsigned MaxSizeWithWaveCount |
| 694 | = ST.getMaxLocalMemSizeWithWaveCount(MaxOccupancy, F); |
| 695 | |
| 696 | // Program is possibly broken by using more local mem than available. |
| 697 | if (CurrentLocalMemUsage > MaxSizeWithWaveCount) |
| 698 | return false; |
| 699 | |
| 700 | LocalMemLimit = MaxSizeWithWaveCount; |
| 701 | |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 702 | LLVM_DEBUG(dbgs() << F.getName() << " uses " << CurrentLocalMemUsage |
| 703 | << " bytes of LDS\n" |
| 704 | << " Rounding size to " << MaxSizeWithWaveCount |
| 705 | << " with a maximum occupancy of " << MaxOccupancy << '\n' |
| 706 | << " and " << (LocalMemLimit - CurrentLocalMemUsage) |
| 707 | << " available for promotion\n"); |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 708 | |
| 709 | return true; |
| 710 | } |
| 711 | |
Matt Arsenault | 8a028bf | 2016-05-16 21:19:59 +0000 | [diff] [blame] | 712 | // FIXME: Should try to pick the most likely to be profitable allocas first. |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 713 | bool AMDGPUPromoteAlloca::handleAlloca(AllocaInst &I, bool SufficientLDS) { |
Matt Arsenault | c5fce69 | 2016-04-28 18:38:48 +0000 | [diff] [blame] | 714 | // Array allocations are probably not worth handling, since an allocation of |
| 715 | // the array type is the canonical form. |
| 716 | if (!I.isStaticAlloca() || I.isArrayAllocation()) |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 717 | return false; |
Matt Arsenault | 19c5488 | 2015-08-26 18:37:13 +0000 | [diff] [blame] | 718 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 719 | IRBuilder<> Builder(&I); |
| 720 | |
| 721 | // First try to replace the alloca with a vector |
| 722 | Type *AllocaTy = I.getAllocatedType(); |
| 723 | |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 724 | LLVM_DEBUG(dbgs() << "Trying to promote " << I << '\n'); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 725 | |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 726 | if (tryPromoteAllocaToVector(&I)) |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 727 | return true; // Promoted to vector. |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 728 | |
Yaxun Liu | 73bf0af | 2018-11-06 21:28:17 +0000 | [diff] [blame] | 729 | if (DisablePromoteAllocaToLDS) |
| 730 | return false; |
| 731 | |
Tom Stellard | 79a1fd7 | 2016-04-14 16:27:07 +0000 | [diff] [blame] | 732 | const Function &ContainingFunction = *I.getParent()->getParent(); |
Matt Arsenault | 5c80618 | 2017-05-02 18:33:18 +0000 | [diff] [blame] | 733 | CallingConv::ID CC = ContainingFunction.getCallingConv(); |
Tom Stellard | 79a1fd7 | 2016-04-14 16:27:07 +0000 | [diff] [blame] | 734 | |
Nicolai Haehnle | bef1ceb | 2016-07-18 09:02:47 +0000 | [diff] [blame] | 735 | // Don't promote the alloca to LDS for shader calling conventions as the work |
| 736 | // item ID intrinsics are not supported for these calling conventions. |
| 737 | // Furthermore not all LDS is available for some of the stages. |
Matt Arsenault | 5c80618 | 2017-05-02 18:33:18 +0000 | [diff] [blame] | 738 | switch (CC) { |
| 739 | case CallingConv::AMDGPU_KERNEL: |
| 740 | case CallingConv::SPIR_KERNEL: |
| 741 | break; |
| 742 | default: |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 743 | LLVM_DEBUG( |
| 744 | dbgs() |
| 745 | << " promote alloca to LDS not supported with calling convention.\n"); |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 746 | return false; |
Matt Arsenault | 5c80618 | 2017-05-02 18:33:18 +0000 | [diff] [blame] | 747 | } |
Nicolai Haehnle | bef1ceb | 2016-07-18 09:02:47 +0000 | [diff] [blame] | 748 | |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 749 | // Not likely to have sufficient local memory for promotion. |
| 750 | if (!SufficientLDS) |
| 751 | return false; |
| 752 | |
Tom Stellard | 5bfbae5 | 2018-07-11 20:59:01 +0000 | [diff] [blame] | 753 | const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(*TM, ContainingFunction); |
Konstantin Zhuravlyov | 1d65026 | 2016-09-06 20:22:28 +0000 | [diff] [blame] | 754 | unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(ContainingFunction).second; |
Tom Stellard | 79a1fd7 | 2016-04-14 16:27:07 +0000 | [diff] [blame] | 755 | |
Matt Arsenault | 8a028bf | 2016-05-16 21:19:59 +0000 | [diff] [blame] | 756 | const DataLayout &DL = Mod->getDataLayout(); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 757 | |
Matt Arsenault | 8a028bf | 2016-05-16 21:19:59 +0000 | [diff] [blame] | 758 | unsigned Align = I.getAlignment(); |
| 759 | if (Align == 0) |
| 760 | Align = DL.getABITypeAlignment(I.getAllocatedType()); |
| 761 | |
| 762 | // FIXME: This computed padding is likely wrong since it depends on inverse |
| 763 | // usage order. |
| 764 | // |
| 765 | // FIXME: It is also possible that if we're allowed to use all of the memory |
| 766 | // could could end up using more than the maximum due to alignment padding. |
| 767 | |
| 768 | uint32_t NewSize = alignTo(CurrentLocalMemUsage, Align); |
| 769 | uint32_t AllocSize = WorkGroupSize * DL.getTypeAllocSize(AllocaTy); |
| 770 | NewSize += AllocSize; |
| 771 | |
| 772 | if (NewSize > LocalMemLimit) { |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 773 | LLVM_DEBUG(dbgs() << " " << AllocSize |
| 774 | << " bytes of local memory not available to promote\n"); |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 775 | return false; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 776 | } |
| 777 | |
Matt Arsenault | 8a028bf | 2016-05-16 21:19:59 +0000 | [diff] [blame] | 778 | CurrentLocalMemUsage = NewSize; |
| 779 | |
Tom Stellard | 5b2927f | 2014-10-31 20:52:04 +0000 | [diff] [blame] | 780 | std::vector<Value*> WorkList; |
| 781 | |
Matt Arsenault | a61cb48 | 2016-05-12 01:58:58 +0000 | [diff] [blame] | 782 | if (!collectUsesWithPtrTypes(&I, &I, WorkList)) { |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 783 | LLVM_DEBUG(dbgs() << " Do not know how to convert all uses\n"); |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 784 | return false; |
Tom Stellard | 5b2927f | 2014-10-31 20:52:04 +0000 | [diff] [blame] | 785 | } |
| 786 | |
Nicola Zaghen | d34e60c | 2018-05-14 12:53:11 +0000 | [diff] [blame] | 787 | LLVM_DEBUG(dbgs() << "Promoting alloca to local memory\n"); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 788 | |
Matt Arsenault | cf84e26 | 2016-02-05 19:47:23 +0000 | [diff] [blame] | 789 | Function *F = I.getParent()->getParent(); |
| 790 | |
Tom Stellard | 79a1fd7 | 2016-04-14 16:27:07 +0000 | [diff] [blame] | 791 | Type *GVTy = ArrayType::get(I.getAllocatedType(), WorkGroupSize); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 792 | GlobalVariable *GV = new GlobalVariable( |
Matt Arsenault | cf84e26 | 2016-02-05 19:47:23 +0000 | [diff] [blame] | 793 | *Mod, GVTy, false, GlobalValue::InternalLinkage, |
| 794 | UndefValue::get(GVTy), |
| 795 | Twine(F->getName()) + Twine('.') + I.getName(), |
| 796 | nullptr, |
| 797 | GlobalVariable::NotThreadLocal, |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 798 | AMDGPUAS::LOCAL_ADDRESS); |
Peter Collingbourne | 96efdd6 | 2016-06-14 21:01:22 +0000 | [diff] [blame] | 799 | GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global); |
Matt Arsenault | cf84e26 | 2016-02-05 19:47:23 +0000 | [diff] [blame] | 800 | GV->setAlignment(I.getAlignment()); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 801 | |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 802 | Value *TCntY, *TCntZ; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 803 | |
Matt Arsenault | e013246 | 2016-01-30 05:19:45 +0000 | [diff] [blame] | 804 | std::tie(TCntY, TCntZ) = getLocalSizeYZ(Builder); |
| 805 | Value *TIdX = getWorkitemID(Builder, 0); |
| 806 | Value *TIdY = getWorkitemID(Builder, 1); |
| 807 | Value *TIdZ = getWorkitemID(Builder, 2); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 808 | |
Matt Arsenault | 853a1fc | 2016-02-02 19:18:48 +0000 | [diff] [blame] | 809 | Value *Tmp0 = Builder.CreateMul(TCntY, TCntZ, "", true, true); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 810 | Tmp0 = Builder.CreateMul(Tmp0, TIdX); |
Matt Arsenault | 853a1fc | 2016-02-02 19:18:48 +0000 | [diff] [blame] | 811 | Value *Tmp1 = Builder.CreateMul(TIdY, TCntZ, "", true, true); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 812 | Value *TID = Builder.CreateAdd(Tmp0, Tmp1); |
| 813 | TID = Builder.CreateAdd(TID, TIdZ); |
| 814 | |
Matt Arsenault | 853a1fc | 2016-02-02 19:18:48 +0000 | [diff] [blame] | 815 | Value *Indices[] = { |
| 816 | Constant::getNullValue(Type::getInt32Ty(Mod->getContext())), |
| 817 | TID |
| 818 | }; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 819 | |
Matt Arsenault | 853a1fc | 2016-02-02 19:18:48 +0000 | [diff] [blame] | 820 | Value *Offset = Builder.CreateInBoundsGEP(GVTy, GV, Indices); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 821 | I.mutateType(Offset->getType()); |
| 822 | I.replaceAllUsesWith(Offset); |
| 823 | I.eraseFromParent(); |
| 824 | |
Matt Arsenault | fb8cdba | 2016-02-02 19:32:35 +0000 | [diff] [blame] | 825 | for (Value *V : WorkList) { |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 826 | CallInst *Call = dyn_cast<CallInst>(V); |
| 827 | if (!Call) { |
Matt Arsenault | 891fccc | 2016-05-18 15:57:21 +0000 | [diff] [blame] | 828 | if (ICmpInst *CI = dyn_cast<ICmpInst>(V)) { |
| 829 | Value *Src0 = CI->getOperand(0); |
| 830 | Type *EltTy = Src0->getType()->getPointerElementType(); |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 831 | PointerType *NewTy = PointerType::get(EltTy, AMDGPUAS::LOCAL_ADDRESS); |
Matt Arsenault | 891fccc | 2016-05-18 15:57:21 +0000 | [diff] [blame] | 832 | |
| 833 | if (isa<ConstantPointerNull>(CI->getOperand(0))) |
| 834 | CI->setOperand(0, ConstantPointerNull::get(NewTy)); |
| 835 | |
| 836 | if (isa<ConstantPointerNull>(CI->getOperand(1))) |
| 837 | CI->setOperand(1, ConstantPointerNull::get(NewTy)); |
| 838 | |
| 839 | continue; |
| 840 | } |
Matt Arsenault | 65f67e4 | 2014-09-15 15:41:44 +0000 | [diff] [blame] | 841 | |
Matt Arsenault | 2402b95 | 2016-12-10 00:52:50 +0000 | [diff] [blame] | 842 | // The operand's value should be corrected on its own and we don't want to |
| 843 | // touch the users. |
Matt Arsenault | 65f67e4 | 2014-09-15 15:41:44 +0000 | [diff] [blame] | 844 | if (isa<AddrSpaceCastInst>(V)) |
| 845 | continue; |
| 846 | |
Matt Arsenault | 891fccc | 2016-05-18 15:57:21 +0000 | [diff] [blame] | 847 | Type *EltTy = V->getType()->getPointerElementType(); |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 848 | PointerType *NewTy = PointerType::get(EltTy, AMDGPUAS::LOCAL_ADDRESS); |
Matt Arsenault | 891fccc | 2016-05-18 15:57:21 +0000 | [diff] [blame] | 849 | |
Matt Arsenault | 65f67e4 | 2014-09-15 15:41:44 +0000 | [diff] [blame] | 850 | // FIXME: It doesn't really make sense to try to do this for all |
| 851 | // instructions. |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 852 | V->mutateType(NewTy); |
Matt Arsenault | 891fccc | 2016-05-18 15:57:21 +0000 | [diff] [blame] | 853 | |
| 854 | // Adjust the types of any constant operands. |
| 855 | if (SelectInst *SI = dyn_cast<SelectInst>(V)) { |
| 856 | if (isa<ConstantPointerNull>(SI->getOperand(1))) |
| 857 | SI->setOperand(1, ConstantPointerNull::get(NewTy)); |
| 858 | |
| 859 | if (isa<ConstantPointerNull>(SI->getOperand(2))) |
| 860 | SI->setOperand(2, ConstantPointerNull::get(NewTy)); |
| 861 | } else if (PHINode *Phi = dyn_cast<PHINode>(V)) { |
| 862 | for (unsigned I = 0, E = Phi->getNumIncomingValues(); I != E; ++I) { |
| 863 | if (isa<ConstantPointerNull>(Phi->getIncomingValue(I))) |
| 864 | Phi->setIncomingValue(I, ConstantPointerNull::get(NewTy)); |
| 865 | } |
| 866 | } |
| 867 | |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 868 | continue; |
| 869 | } |
| 870 | |
Matt Arsenault | 2e08e18 | 2016-07-18 18:34:48 +0000 | [diff] [blame] | 871 | IntrinsicInst *Intr = cast<IntrinsicInst>(Call); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 872 | Builder.SetInsertPoint(Intr); |
| 873 | switch (Intr->getIntrinsicID()) { |
| 874 | case Intrinsic::lifetime_start: |
| 875 | case Intrinsic::lifetime_end: |
| 876 | // These intrinsics are for address space 0 only |
| 877 | Intr->eraseFromParent(); |
| 878 | continue; |
| 879 | case Intrinsic::memcpy: { |
| 880 | MemCpyInst *MemCpy = cast<MemCpyInst>(Intr); |
Daniel Neilson | a60f462 | 2018-02-09 21:56:15 +0000 | [diff] [blame] | 881 | Builder.CreateMemCpy(MemCpy->getRawDest(), MemCpy->getDestAlignment(), |
| 882 | MemCpy->getRawSource(), MemCpy->getSourceAlignment(), |
| 883 | MemCpy->getLength(), MemCpy->isVolatile()); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 884 | Intr->eraseFromParent(); |
| 885 | continue; |
| 886 | } |
Matt Arsenault | 7e747f1 | 2016-02-02 20:28:10 +0000 | [diff] [blame] | 887 | case Intrinsic::memmove: { |
| 888 | MemMoveInst *MemMove = cast<MemMoveInst>(Intr); |
Daniel Neilson | a60f462 | 2018-02-09 21:56:15 +0000 | [diff] [blame] | 889 | Builder.CreateMemMove(MemMove->getRawDest(), MemMove->getDestAlignment(), |
| 890 | MemMove->getRawSource(), MemMove->getSourceAlignment(), |
| 891 | MemMove->getLength(), MemMove->isVolatile()); |
Matt Arsenault | 7e747f1 | 2016-02-02 20:28:10 +0000 | [diff] [blame] | 892 | Intr->eraseFromParent(); |
| 893 | continue; |
| 894 | } |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 895 | case Intrinsic::memset: { |
| 896 | MemSetInst *MemSet = cast<MemSetInst>(Intr); |
| 897 | Builder.CreateMemSet(MemSet->getRawDest(), MemSet->getValue(), |
Daniel Neilson | a60f462 | 2018-02-09 21:56:15 +0000 | [diff] [blame] | 898 | MemSet->getLength(), MemSet->getDestAlignment(), |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 899 | MemSet->isVolatile()); |
| 900 | Intr->eraseFromParent(); |
| 901 | continue; |
| 902 | } |
Matt Arsenault | 0b783ef0 | 2016-01-22 19:47:54 +0000 | [diff] [blame] | 903 | case Intrinsic::invariant_start: |
| 904 | case Intrinsic::invariant_end: |
Piotr Padlewski | 5dde809 | 2018-05-03 11:03:01 +0000 | [diff] [blame] | 905 | case Intrinsic::launder_invariant_group: |
Piotr Padlewski | 5b3db45 | 2018-07-02 04:49:30 +0000 | [diff] [blame] | 906 | case Intrinsic::strip_invariant_group: |
Matt Arsenault | 0b783ef0 | 2016-01-22 19:47:54 +0000 | [diff] [blame] | 907 | Intr->eraseFromParent(); |
| 908 | // FIXME: I think the invariant marker should still theoretically apply, |
| 909 | // but the intrinsics need to be changed to accept pointers with any |
| 910 | // address space. |
| 911 | continue; |
Matt Arsenault | 7e747f1 | 2016-02-02 20:28:10 +0000 | [diff] [blame] | 912 | case Intrinsic::objectsize: { |
| 913 | Value *Src = Intr->getOperand(0); |
| 914 | Type *SrcTy = Src->getType()->getPointerElementType(); |
| 915 | Function *ObjectSize = Intrinsic::getDeclaration(Mod, |
| 916 | Intrinsic::objectsize, |
Matt Arsenault | 0da6350 | 2018-08-31 05:49:54 +0000 | [diff] [blame] | 917 | { Intr->getType(), PointerType::get(SrcTy, AMDGPUAS::LOCAL_ADDRESS) } |
Matt Arsenault | 7e747f1 | 2016-02-02 20:28:10 +0000 | [diff] [blame] | 918 | ); |
| 919 | |
George Burgess IV | 56c7e88 | 2017-03-21 20:08:59 +0000 | [diff] [blame] | 920 | CallInst *NewCall = Builder.CreateCall( |
Erik Pilkington | 600e9de | 2019-01-30 20:34:35 +0000 | [diff] [blame] | 921 | ObjectSize, |
| 922 | {Src, Intr->getOperand(1), Intr->getOperand(2), Intr->getOperand(3)}); |
Matt Arsenault | 7e747f1 | 2016-02-02 20:28:10 +0000 | [diff] [blame] | 923 | Intr->replaceAllUsesWith(NewCall); |
| 924 | Intr->eraseFromParent(); |
| 925 | continue; |
| 926 | } |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 927 | default: |
Matthias Braun | 8c209aa | 2017-01-28 02:02:38 +0000 | [diff] [blame] | 928 | Intr->print(errs()); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 929 | llvm_unreachable("Don't know how to promote alloca intrinsic use."); |
| 930 | } |
| 931 | } |
Changpeng Fang | 1dbace1 | 2017-05-23 20:25:41 +0000 | [diff] [blame] | 932 | return true; |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 933 | } |
| 934 | |
Francis Visoiu Mistrih | 8b61764 | 2017-05-18 17:21:13 +0000 | [diff] [blame] | 935 | FunctionPass *llvm::createAMDGPUPromoteAlloca() { |
| 936 | return new AMDGPUPromoteAlloca(); |
Tom Stellard | 880a80a | 2014-06-17 16:53:14 +0000 | [diff] [blame] | 937 | } |