LLVM 17.0.0git
NVPTXLowerArgs.cpp
Go to the documentation of this file.
1//===-- NVPTXLowerArgs.cpp - Lower arguments ------------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9//
10// Arguments to kernel and device functions are passed via param space,
11// which imposes certain restrictions:
12// http://docs.nvidia.com/cuda/parallel-thread-execution/#state-spaces
13//
14// Kernel parameters are read-only and accessible only via ld.param
15// instruction, directly or via a pointer. Pointers to kernel
16// arguments can't be converted to generic address space.
17//
18// Device function parameters are directly accessible via
19// ld.param/st.param, but taking the address of one returns a pointer
20// to a copy created in local space which *can't* be used with
21// ld.param/st.param.
22//
23// Copying a byval struct into local memory in IR allows us to enforce
24// the param space restrictions, gives the rest of IR a pointer w/o
25// param space restrictions, and gives us an opportunity to eliminate
26// the copy.
27//
28// Pointer arguments to kernel functions need more work to be lowered:
29//
30// 1. Convert non-byval pointer arguments of CUDA kernels to pointers in the
31// global address space. This allows later optimizations to emit
32// ld.global.*/st.global.* for accessing these pointer arguments. For
33// example,
34//
35// define void @foo(float* %input) {
36// %v = load float, float* %input, align 4
37// ...
38// }
39//
40// becomes
41//
42// define void @foo(float* %input) {
43// %input2 = addrspacecast float* %input to float addrspace(1)*
44// %input3 = addrspacecast float addrspace(1)* %input2 to float*
45// %v = load float, float* %input3, align 4
46// ...
47// }
48//
49// Later, NVPTXInferAddressSpaces will optimize it to
50//
51// define void @foo(float* %input) {
52// %input2 = addrspacecast float* %input to float addrspace(1)*
53// %v = load float, float addrspace(1)* %input2, align 4
54// ...
55// }
56//
57// 2. Convert pointers in a byval kernel parameter to pointers in the global
58// address space. As #2, it allows NVPTX to emit more ld/st.global. E.g.,
59//
60// struct S {
61// int *x;
62// int *y;
63// };
64// __global__ void foo(S s) {
65// int *b = s.y;
66// // use b
67// }
68//
69// "b" points to the global address space. In the IR level,
70//
71// define void @foo({i32*, i32*}* byval %input) {
72// %b_ptr = getelementptr {i32*, i32*}, {i32*, i32*}* %input, i64 0, i32 1
73// %b = load i32*, i32** %b_ptr
74// ; use %b
75// }
76//
77// becomes
78//
79// define void @foo({i32*, i32*}* byval %input) {
80// %b_ptr = getelementptr {i32*, i32*}, {i32*, i32*}* %input, i64 0, i32 1
81// %b = load i32*, i32** %b_ptr
82// %b_global = addrspacecast i32* %b to i32 addrspace(1)*
83// %b_generic = addrspacecast i32 addrspace(1)* %b_global to i32*
84// ; use %b_generic
85// }
86//
87// TODO: merge this pass with NVPTXInferAddressSpaces so that other passes don't
88// cancel the addrspacecast pair this pass emits.
89//===----------------------------------------------------------------------===//
90
92#include "NVPTX.h"
93#include "NVPTXTargetMachine.h"
94#include "NVPTXUtilities.h"
97#include "llvm/IR/Function.h"
99#include "llvm/IR/Module.h"
100#include "llvm/IR/Type.h"
102#include "llvm/Pass.h"
103#include <numeric>
104#include <queue>
105
106#define DEBUG_TYPE "nvptx-lower-args"
107
108using namespace llvm;
109
110namespace llvm {
112}
113
114namespace {
115class NVPTXLowerArgs : public FunctionPass {
116 bool runOnFunction(Function &F) override;
117
118 bool runOnKernelFunction(const NVPTXTargetMachine &TM, Function &F);
119 bool runOnDeviceFunction(const NVPTXTargetMachine &TM, Function &F);
120
121 // handle byval parameters
122 void handleByValParam(const NVPTXTargetMachine &TM, Argument *Arg);
123 // Knowing Ptr must point to the global address space, this function
124 // addrspacecasts Ptr to global and then back to generic. This allows
125 // NVPTXInferAddressSpaces to fold the global-to-generic cast into
126 // loads/stores that appear later.
127 void markPointerAsGlobal(Value *Ptr);
128
129public:
130 static char ID; // Pass identification, replacement for typeid
131 NVPTXLowerArgs() : FunctionPass(ID) {}
132 StringRef getPassName() const override {
133 return "Lower pointer arguments of CUDA kernels";
134 }
135 void getAnalysisUsage(AnalysisUsage &AU) const override {
137 }
138};
139} // namespace
140
141char NVPTXLowerArgs::ID = 1;
142
143INITIALIZE_PASS_BEGIN(NVPTXLowerArgs, "nvptx-lower-args",
144 "Lower arguments (NVPTX)", false, false)
146INITIALIZE_PASS_END(NVPTXLowerArgs, "nvptx-lower-args",
148
149// =============================================================================
150// If the function had a byval struct ptr arg, say foo(%struct.x* byval %d),
151// and we can't guarantee that the only accesses are loads,
152// then add the following instructions to the first basic block:
153//
154// %temp = alloca %struct.x, align 8
155// %tempd = addrspacecast %struct.x* %d to %struct.x addrspace(101)*
156// %tv = load %struct.x addrspace(101)* %tempd
157// store %struct.x %tv, %struct.x* %temp, align 8
158//
159// The above code allocates some space in the stack and copies the incoming
160// struct from param space to local space.
161// Then replace all occurrences of %d by %temp.
162//
163// In case we know that all users are GEPs or Loads, replace them with the same
164// ones in parameter AS, so we can access them using ld.param.
165// =============================================================================
166
167// Replaces the \p OldUser instruction with the same in parameter AS.
168// Only Load and GEP are supported.
169static void convertToParamAS(Value *OldUser, Value *Param) {
170 Instruction *I = dyn_cast<Instruction>(OldUser);
171 assert(I && "OldUser must be an instruction");
172 struct IP {
173 Instruction *OldInstruction;
174 Value *NewParam;
175 };
176 SmallVector<IP> ItemsToConvert = {{I, Param}};
177 SmallVector<Instruction *> InstructionsToDelete;
178
179 auto CloneInstInParamAS = [](const IP &I) -> Value * {
180 if (auto *LI = dyn_cast<LoadInst>(I.OldInstruction)) {
181 LI->setOperand(0, I.NewParam);
182 return LI;
183 }
184 if (auto *GEP = dyn_cast<GetElementPtrInst>(I.OldInstruction)) {
185 SmallVector<Value *, 4> Indices(GEP->indices());
186 auto *NewGEP = GetElementPtrInst::Create(GEP->getSourceElementType(),
187 I.NewParam, Indices,
188 GEP->getName(), GEP);
189 NewGEP->setIsInBounds(GEP->isInBounds());
190 return NewGEP;
191 }
192 if (auto *BC = dyn_cast<BitCastInst>(I.OldInstruction)) {
193 auto *NewBCType = PointerType::getWithSamePointeeType(
194 cast<PointerType>(BC->getType()), ADDRESS_SPACE_PARAM);
195 return BitCastInst::Create(BC->getOpcode(), I.NewParam, NewBCType,
196 BC->getName(), BC);
197 }
198 if (auto *ASC = dyn_cast<AddrSpaceCastInst>(I.OldInstruction)) {
199 assert(ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM);
200 (void)ASC;
201 // Just pass through the argument, the old ASC is no longer needed.
202 return I.NewParam;
203 }
204 llvm_unreachable("Unsupported instruction");
205 };
206
207 while (!ItemsToConvert.empty()) {
208 IP I = ItemsToConvert.pop_back_val();
209 Value *NewInst = CloneInstInParamAS(I);
210
211 if (NewInst && NewInst != I.OldInstruction) {
212 // We've created a new instruction. Queue users of the old instruction to
213 // be converted and the instruction itself to be deleted. We can't delete
214 // the old instruction yet, because it's still in use by a load somewhere.
215 for (Value *V : I.OldInstruction->users())
216 ItemsToConvert.push_back({cast<Instruction>(V), NewInst});
217
218 InstructionsToDelete.push_back(I.OldInstruction);
219 }
220 }
221
222 // Now we know that all argument loads are using addresses in parameter space
223 // and we can finally remove the old instructions in generic AS. Instructions
224 // scheduled for removal should be processed in reverse order so the ones
225 // closest to the load are deleted first. Otherwise they may still be in use.
226 // E.g if we have Value = Load(BitCast(GEP(arg))), InstructionsToDelete will
227 // have {GEP,BitCast}. GEP can't be deleted first, because it's still used by
228 // the BitCast.
229 for (Instruction *I : llvm::reverse(InstructionsToDelete))
230 I->eraseFromParent();
231}
232
233// Adjust alignment of arguments passed byval in .param address space. We can
234// increase alignment of such arguments in a way that ensures that we can
235// effectively vectorize their loads. We should also traverse all loads from
236// byval pointer and adjust their alignment, if those were using known offset.
237// Such alignment changes must be conformed with parameter store and load in
238// NVPTXTargetLowering::LowerCall.
239static void adjustByValArgAlignment(Argument *Arg, Value *ArgInParamAS,
240 const NVPTXTargetLowering *TLI) {
241 Function *Func = Arg->getParent();
242 Type *StructType = Arg->getParamByValType();
243 const DataLayout DL(Func->getParent());
244
245 uint64_t NewArgAlign =
247 uint64_t CurArgAlign =
248 Arg->getAttribute(Attribute::Alignment).getValueAsInt();
249
250 if (CurArgAlign >= NewArgAlign)
251 return;
252
253 LLVM_DEBUG(dbgs() << "Try to use alignment " << NewArgAlign << " instead of "
254 << CurArgAlign << " for " << *Arg << '\n');
255
256 auto NewAlignAttr =
257 Attribute::get(Func->getContext(), Attribute::Alignment, NewArgAlign);
258 Arg->removeAttr(Attribute::Alignment);
259 Arg->addAttr(NewAlignAttr);
260
261 struct Load {
262 LoadInst *Inst;
264 };
265
266 struct LoadContext {
267 Value *InitialVal;
269 };
270
271 SmallVector<Load> Loads;
272 std::queue<LoadContext> Worklist;
273 Worklist.push({ArgInParamAS, 0});
274
275 while (!Worklist.empty()) {
276 LoadContext Ctx = Worklist.front();
277 Worklist.pop();
278
279 for (User *CurUser : Ctx.InitialVal->users()) {
280 if (auto *I = dyn_cast<LoadInst>(CurUser)) {
281 Loads.push_back({I, Ctx.Offset});
282 continue;
283 }
284
285 if (auto *I = dyn_cast<BitCastInst>(CurUser)) {
286 Worklist.push({I, Ctx.Offset});
287 continue;
288 }
289
290 if (auto *I = dyn_cast<GetElementPtrInst>(CurUser)) {
291 APInt OffsetAccumulated =
292 APInt::getZero(DL.getIndexSizeInBits(ADDRESS_SPACE_PARAM));
293
294 if (!I->accumulateConstantOffset(DL, OffsetAccumulated))
295 continue;
296
297 uint64_t OffsetLimit = -1;
298 uint64_t Offset = OffsetAccumulated.getLimitedValue(OffsetLimit);
299 assert(Offset != OffsetLimit && "Expect Offset less than UINT64_MAX");
300
301 Worklist.push({I, Ctx.Offset + Offset});
302 continue;
303 }
304
305 llvm_unreachable("All users must be one of: load, "
306 "bitcast, getelementptr.");
307 }
308 }
309
310 for (Load &CurLoad : Loads) {
311 Align NewLoadAlign(std::gcd(NewArgAlign, CurLoad.Offset));
312 Align CurLoadAlign(CurLoad.Inst->getAlign());
313 CurLoad.Inst->setAlignment(std::max(NewLoadAlign, CurLoadAlign));
314 }
315}
316
317void NVPTXLowerArgs::handleByValParam(const NVPTXTargetMachine &TM,
318 Argument *Arg) {
319 Function *Func = Arg->getParent();
320 Instruction *FirstInst = &(Func->getEntryBlock().front());
321 Type *StructType = Arg->getParamByValType();
322 assert(StructType && "Missing byval type");
323
324 auto IsALoadChain = [&](Value *Start) {
325 SmallVector<Value *, 16> ValuesToCheck = {Start};
326 auto IsALoadChainInstr = [](Value *V) -> bool {
327 if (isa<GetElementPtrInst>(V) || isa<BitCastInst>(V) || isa<LoadInst>(V))
328 return true;
329 // ASC to param space are OK, too -- we'll just strip them.
330 if (auto *ASC = dyn_cast<AddrSpaceCastInst>(V)) {
331 if (ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM)
332 return true;
333 }
334 return false;
335 };
336
337 while (!ValuesToCheck.empty()) {
338 Value *V = ValuesToCheck.pop_back_val();
339 if (!IsALoadChainInstr(V)) {
340 LLVM_DEBUG(dbgs() << "Need a copy of " << *Arg << " because of " << *V
341 << "\n");
342 (void)Arg;
343 return false;
344 }
345 if (!isa<LoadInst>(V))
346 llvm::append_range(ValuesToCheck, V->users());
347 }
348 return true;
349 };
350
351 if (llvm::all_of(Arg->users(), IsALoadChain)) {
352 // Convert all loads and intermediate operations to use parameter AS and
353 // skip creation of a local copy of the argument.
354 SmallVector<User *, 16> UsersToUpdate(Arg->users());
355 Value *ArgInParamAS = new AddrSpaceCastInst(
356 Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(),
357 FirstInst);
358 for (Value *V : UsersToUpdate)
359 convertToParamAS(V, ArgInParamAS);
360 LLVM_DEBUG(dbgs() << "No need to copy " << *Arg << "\n");
361
362 const auto *TLI =
363 cast<NVPTXTargetLowering>(TM.getSubtargetImpl()->getTargetLowering());
364
365 adjustByValArgAlignment(Arg, ArgInParamAS, TLI);
366
367 return;
368 }
369
370 // Otherwise we have to create a temporary copy.
371 const DataLayout &DL = Func->getParent()->getDataLayout();
372 unsigned AS = DL.getAllocaAddrSpace();
373 AllocaInst *AllocA = new AllocaInst(StructType, AS, Arg->getName(), FirstInst);
374 // Set the alignment to alignment of the byval parameter. This is because,
375 // later load/stores assume that alignment, and we are going to replace
376 // the use of the byval parameter with this alloca instruction.
377 AllocA->setAlignment(Func->getParamAlign(Arg->getArgNo())
378 .value_or(DL.getPrefTypeAlign(StructType)));
379 Arg->replaceAllUsesWith(AllocA);
380
381 Value *ArgInParam = new AddrSpaceCastInst(
382 Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(),
383 FirstInst);
384 // Be sure to propagate alignment to this load; LLVM doesn't know that NVPTX
385 // addrspacecast preserves alignment. Since params are constant, this load is
386 // definitely not volatile.
387 LoadInst *LI =
388 new LoadInst(StructType, ArgInParam, Arg->getName(),
389 /*isVolatile=*/false, AllocA->getAlign(), FirstInst);
390 new StoreInst(LI, AllocA, FirstInst);
391}
392
393void NVPTXLowerArgs::markPointerAsGlobal(Value *Ptr) {
394 if (Ptr->getType()->getPointerAddressSpace() != ADDRESS_SPACE_GENERIC)
395 return;
396
397 // Deciding where to emit the addrspacecast pair.
398 BasicBlock::iterator InsertPt;
399 if (Argument *Arg = dyn_cast<Argument>(Ptr)) {
400 // Insert at the functon entry if Ptr is an argument.
401 InsertPt = Arg->getParent()->getEntryBlock().begin();
402 } else {
403 // Insert right after Ptr if Ptr is an instruction.
404 InsertPt = ++cast<Instruction>(Ptr)->getIterator();
405 assert(InsertPt != InsertPt->getParent()->end() &&
406 "We don't call this function with Ptr being a terminator.");
407 }
408
409 Instruction *PtrInGlobal = new AddrSpaceCastInst(
410 Ptr,
411 PointerType::getWithSamePointeeType(cast<PointerType>(Ptr->getType()),
413 Ptr->getName(), &*InsertPt);
414 Value *PtrInGeneric = new AddrSpaceCastInst(PtrInGlobal, Ptr->getType(),
415 Ptr->getName(), &*InsertPt);
416 // Replace with PtrInGeneric all uses of Ptr except PtrInGlobal.
417 Ptr->replaceAllUsesWith(PtrInGeneric);
418 PtrInGlobal->setOperand(0, Ptr);
419}
420
421// =============================================================================
422// Main function for this pass.
423// =============================================================================
424bool NVPTXLowerArgs::runOnKernelFunction(const NVPTXTargetMachine &TM,
425 Function &F) {
426 // Copying of byval aggregates + SROA may result in pointers being loaded as
427 // integers, followed by intotoptr. We may want to mark those as global, too,
428 // but only if the loaded integer is used exclusively for conversion to a
429 // pointer with inttoptr.
430 auto HandleIntToPtr = [this](Value &V) {
431 if (llvm::all_of(V.users(), [](User *U) { return isa<IntToPtrInst>(U); })) {
432 SmallVector<User *, 16> UsersToUpdate(V.users());
433 llvm::for_each(UsersToUpdate, [&](User *U) { markPointerAsGlobal(U); });
434 }
435 };
436 if (TM.getDrvInterface() == NVPTX::CUDA) {
437 // Mark pointers in byval structs as global.
438 for (auto &B : F) {
439 for (auto &I : B) {
440 if (LoadInst *LI = dyn_cast<LoadInst>(&I)) {
441 if (LI->getType()->isPointerTy() || LI->getType()->isIntegerTy()) {
442 Value *UO = getUnderlyingObject(LI->getPointerOperand());
443 if (Argument *Arg = dyn_cast<Argument>(UO)) {
444 if (Arg->hasByValAttr()) {
445 // LI is a load from a pointer within a byval kernel parameter.
446 if (LI->getType()->isPointerTy())
447 markPointerAsGlobal(LI);
448 else
449 HandleIntToPtr(*LI);
450 }
451 }
452 }
453 }
454 }
455 }
456 }
457
458 LLVM_DEBUG(dbgs() << "Lowering kernel args of " << F.getName() << "\n");
459 for (Argument &Arg : F.args()) {
460 if (Arg.getType()->isPointerTy()) {
461 if (Arg.hasByValAttr())
462 handleByValParam(TM, &Arg);
463 else if (TM.getDrvInterface() == NVPTX::CUDA)
464 markPointerAsGlobal(&Arg);
465 } else if (Arg.getType()->isIntegerTy() &&
466 TM.getDrvInterface() == NVPTX::CUDA) {
467 HandleIntToPtr(Arg);
468 }
469 }
470 return true;
471}
472
473// Device functions only need to copy byval args into local memory.
474bool NVPTXLowerArgs::runOnDeviceFunction(const NVPTXTargetMachine &TM,
475 Function &F) {
476 LLVM_DEBUG(dbgs() << "Lowering function args of " << F.getName() << "\n");
477 for (Argument &Arg : F.args())
478 if (Arg.getType()->isPointerTy() && Arg.hasByValAttr())
479 handleByValParam(TM, &Arg);
480 return true;
481}
482
483bool NVPTXLowerArgs::runOnFunction(Function &F) {
484 auto &TM = getAnalysis<TargetPassConfig>().getTM<NVPTXTargetMachine>();
485
486 return isKernelFunction(F) ? runOnKernelFunction(TM, F)
487 : runOnDeviceFunction(TM, F);
488}
489
490FunctionPass *llvm::createNVPTXLowerArgsPass() { return new NVPTXLowerArgs(); }
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
amdgpu Simplify well known AMD library false FunctionCallee Value * Arg
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
#define LLVM_DEBUG(X)
Definition: Debug.h:101
Hexagon Common GEP
#define F(x, y, z)
Definition: MD5.cpp:55
#define I(x, y, z)
Definition: MD5.cpp:58
Module.h This file contains the declarations for the Module class.
nvptx lower Lower arguments(NVPTX)"
nvptx lower args
static void adjustByValArgAlignment(Argument *Arg, Value *ArgInParamAS, const NVPTXTargetLowering *TLI)
nvptx lower Lower static false void convertToParamAS(Value *OldUser, Value *Param)
const char LLVMTargetMachineRef TM
#define INITIALIZE_PASS_DEPENDENCY(depName)
Definition: PassSupport.h:55
#define INITIALIZE_PASS_END(passName, arg, name, cfg, analysis)
Definition: PassSupport.h:59
#define INITIALIZE_PASS_BEGIN(passName, arg, name, cfg, analysis)
Definition: PassSupport.h:52
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
Target-Independent Code Generator Pass Configuration Options pass.
Class for arbitrary precision integers.
Definition: APInt.h:75
uint64_t getLimitedValue(uint64_t Limit=UINT64_MAX) const
If this value is smaller than the specified limit, return it, otherwise return the limit value.
Definition: APInt.h:463
static APInt getZero(unsigned numBits)
Get the '0' value for the specified bit-width.
Definition: APInt.h:177
This class represents a conversion between pointers from one address space to another.
an instruction to allocate memory on the stack
Definition: Instructions.h:58
Align getAlign() const
Return the alignment of the memory that is being allocated by the instruction.
Definition: Instructions.h:125
void setAlignment(Align Align)
Definition: Instructions.h:129
Represent the analysis usage information of a pass.
AnalysisUsage & addRequired()
This class represents an incoming formal argument to a Function.
Definition: Argument.h:28
static Attribute get(LLVMContext &Context, AttrKind Kind, uint64_t Val=0)
Return a uniquified Attribute object.
Definition: Attributes.cpp:91
InstListType::iterator iterator
Instruction iterators...
Definition: BasicBlock.h:87
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:114
FunctionPass class - This class is used to implement most global optimizations.
Definition: Pass.h:308
virtual bool runOnFunction(Function &F)=0
runOnFunction - Virtual method overriden by subclasses to do the per-function processing of the pass.
static GetElementPtrInst * Create(Type *PointeeType, Value *Ptr, ArrayRef< Value * > IdxList, const Twine &NameStr="", Instruction *InsertBefore=nullptr)
Definition: Instructions.h:966
An instruction for reading from memory.
Definition: Instructions.h:177
Align getFunctionParamOptimizedAlign(const Function *F, Type *ArgTy, const DataLayout &DL) const
getFunctionParamOptimizedAlign - since function arguments are passed via .param space,...
PassRegistry - This class manages the registration and intitialization of the pass subsystem as appli...
Definition: PassRegistry.h:38
virtual void getAnalysisUsage(AnalysisUsage &) const
getAnalysisUsage - This function should be overriden by passes that need analysis information to do t...
Definition: Pass.cpp:98
virtual StringRef getPassName() const
getPassName - Return a nice clean name for a pass.
Definition: Pass.cpp:81
bool empty() const
Definition: SmallVector.h:94
void push_back(const T &Elt)
Definition: SmallVector.h:416
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1200
An instruction for storing to memory.
Definition: Instructions.h:301
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:50
Class to represent struct types.
Definition: DerivedTypes.h:213
Target-Independent Code Generator Pass Configuration Options.
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
void setOperand(unsigned i, Value *Val)
Definition: User.h:174
LLVM Value Representation.
Definition: Value.h:74
iterator_range< user_iterator > users()
Definition: Value.h:421
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Definition: CallingConv.h:24
@ CUDA
Definition: NVPTX.h:73
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:18
@ Offset
Definition: DWP.cpp:406
UnaryFunction for_each(R &&Range, UnaryFunction F)
Provide wrappers to std::for_each which take ranges instead of having to pass begin/end explicitly.
Definition: STLExtras.h:1728
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
Definition: STLExtras.h:1735
@ ADDRESS_SPACE_GENERIC
Definition: NVPTXBaseInfo.h:22
@ ADDRESS_SPACE_GLOBAL
Definition: NVPTXBaseInfo.h:23
@ ADDRESS_SPACE_PARAM
Definition: NVPTXBaseInfo.h:29
void append_range(Container &C, Range &&R)
Wrapper function to append a range to a container.
Definition: STLExtras.h:2014
const Value * getUnderlyingObject(const Value *V, unsigned MaxLookup=6)
This method strips off any GEP address adjustments and pointer casts from the specified value,...
void initializeNVPTXLowerArgsPass(PassRegistry &)
FunctionPass * createNVPTXLowerArgsPass()
auto reverse(ContainerTy &&C)
Definition: STLExtras.h:484
raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition: Debug.cpp:163
bool isKernelFunction(const Function &F)
This struct is a compact representation of a valid (non-zero power of two) alignment.
Definition: Alignment.h:39
uint64_t value() const
This is a hole in the type system and should not be abused.
Definition: Alignment.h:85