LLVM  15.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"
96 #include "llvm/IR/Function.h"
97 #include "llvm/IR/Instructions.h"
98 #include "llvm/IR/Module.h"
99 #include "llvm/IR/Type.h"
100 #include "llvm/Pass.h"
101 #include <queue>
102 
103 #define DEBUG_TYPE "nvptx-lower-args"
104 
105 using namespace llvm;
106 
107 namespace llvm {
109 }
110 
111 namespace {
112 class NVPTXLowerArgs : public FunctionPass {
113  bool runOnFunction(Function &F) override;
114 
115  bool runOnKernelFunction(Function &F);
116  bool runOnDeviceFunction(Function &F);
117 
118  // handle byval parameters
119  void handleByValParam(Argument *Arg);
120  // Knowing Ptr must point to the global address space, this function
121  // addrspacecasts Ptr to global and then back to generic. This allows
122  // NVPTXInferAddressSpaces to fold the global-to-generic cast into
123  // loads/stores that appear later.
124  void markPointerAsGlobal(Value *Ptr);
125 
126 public:
127  static char ID; // Pass identification, replacement for typeid
128  NVPTXLowerArgs(const NVPTXTargetMachine *TM = nullptr)
129  : FunctionPass(ID), TM(TM) {}
130  StringRef getPassName() const override {
131  return "Lower pointer arguments of CUDA kernels";
132  }
133 
134 private:
135  const NVPTXTargetMachine *TM;
136 };
137 } // namespace
138 
139 char NVPTXLowerArgs::ID = 1;
140 
141 INITIALIZE_PASS(NVPTXLowerArgs, "nvptx-lower-args",
142  "Lower arguments (NVPTX)", false, false)
143 
144 // =============================================================================
145 // If the function had a byval struct ptr arg, say foo(%struct.x* byval %d),
146 // and we can't guarantee that the only accesses are loads,
147 // then add the following instructions to the first basic block:
148 //
149 // %temp = alloca %struct.x, align 8
150 // %tempd = addrspacecast %struct.x* %d to %struct.x addrspace(101)*
151 // %tv = load %struct.x addrspace(101)* %tempd
152 // store %struct.x %tv, %struct.x* %temp, align 8
153 //
154 // The above code allocates some space in the stack and copies the incoming
155 // struct from param space to local space.
156 // Then replace all occurrences of %d by %temp.
157 //
158 // In case we know that all users are GEPs or Loads, replace them with the same
159 // ones in parameter AS, so we can access them using ld.param.
160 // =============================================================================
161 
162 // Replaces the \p OldUser instruction with the same in parameter AS.
163 // Only Load and GEP are supported.
164 static void convertToParamAS(Value *OldUser, Value *Param) {
165  Instruction *I = dyn_cast<Instruction>(OldUser);
166  assert(I && "OldUser must be an instruction");
167  struct IP {
170  };
173 
174  auto CloneInstInParamAS = [](const IP &I) -> Value * {
175  if (auto *LI = dyn_cast<LoadInst>(I.OldInstruction)) {
176  LI->setOperand(0, I.NewParam);
177  return LI;
178  }
179  if (auto *GEP = dyn_cast<GetElementPtrInst>(I.OldInstruction)) {
180  SmallVector<Value *, 4> Indices(GEP->indices());
181  auto *NewGEP = GetElementPtrInst::Create(GEP->getSourceElementType(),
182  I.NewParam, Indices,
183  GEP->getName(), GEP);
184  NewGEP->setIsInBounds(GEP->isInBounds());
185  return NewGEP;
186  }
187  if (auto *BC = dyn_cast<BitCastInst>(I.OldInstruction)) {
188  auto *NewBCType = PointerType::getWithSamePointeeType(
189  cast<PointerType>(BC->getType()), ADDRESS_SPACE_PARAM);
190  return BitCastInst::Create(BC->getOpcode(), I.NewParam, NewBCType,
191  BC->getName(), BC);
192  }
193  if (auto *ASC = dyn_cast<AddrSpaceCastInst>(I.OldInstruction)) {
194  assert(ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM);
195  (void)ASC;
196  // Just pass through the argument, the old ASC is no longer needed.
197  return I.NewParam;
198  }
199  llvm_unreachable("Unsupported instruction");
200  };
201 
202  while (!ItemsToConvert.empty()) {
203  IP I = ItemsToConvert.pop_back_val();
204  Value *NewInst = CloneInstInParamAS(I);
205 
206  if (NewInst && NewInst != I.OldInstruction) {
207  // We've created a new instruction. Queue users of the old instruction to
208  // be converted and the instruction itself to be deleted. We can't delete
209  // the old instruction yet, because it's still in use by a load somewhere.
210  for (Value *V : I.OldInstruction->users())
211  ItemsToConvert.push_back({cast<Instruction>(V), NewInst});
212 
213  InstructionsToDelete.push_back(I.OldInstruction);
214  }
215  }
216 
217  // Now we know that all argument loads are using addresses in parameter space
218  // and we can finally remove the old instructions in generic AS. Instructions
219  // scheduled for removal should be processed in reverse order so the ones
220  // closest to the load are deleted first. Otherwise they may still be in use.
221  // E.g if we have Value = Load(BitCast(GEP(arg))), InstructionsToDelete will
222  // have {GEP,BitCast}. GEP can't be deleted first, because it's still used by
223  // the BitCast.
225  I->eraseFromParent();
226 }
227 
228 // Adjust alignment of arguments passed byval in .param address space. We can
229 // increase alignment of such arguments in a way that ensures that we can
230 // effectively vectorize their loads. We should also traverse all loads from
231 // byval pointer and adjust their alignment, if those were using known offset.
232 // Such alignment changes must be conformed with parameter store and load in
233 // NVPTXTargetLowering::LowerCall.
234 static void adjustByValArgAlignment(Argument *Arg, Value *ArgInParamAS,
235  const NVPTXTargetLowering *TLI) {
236  Function *Func = Arg->getParent();
237  Type *StructType = Arg->getParamByValType();
238  const DataLayout DL(Func->getParent());
239 
240  uint64_t NewArgAlign =
242  uint64_t CurArgAlign =
243  Arg->getAttribute(Attribute::Alignment).getValueAsInt();
244 
245  if (CurArgAlign >= NewArgAlign)
246  return;
247 
248  LLVM_DEBUG(dbgs() << "Try to use alignment " << NewArgAlign << " instead of "
249  << CurArgAlign << " for " << *Arg << '\n');
250 
251  auto NewAlignAttr =
252  Attribute::get(Func->getContext(), Attribute::Alignment, NewArgAlign);
253  Arg->removeAttr(Attribute::Alignment);
254  Arg->addAttr(NewAlignAttr);
255 
256  struct Load {
257  LoadInst *Inst;
258  uint64_t Offset;
259  };
260 
261  struct LoadContext {
262  Value *InitialVal;
263  uint64_t Offset;
264  };
265 
266  SmallVector<Load> Loads;
267  std::queue<LoadContext> Worklist;
268  Worklist.push({ArgInParamAS, 0});
269 
270  while (!Worklist.empty()) {
271  LoadContext Ctx = Worklist.front();
272  Worklist.pop();
273 
274  for (User *CurUser : Ctx.InitialVal->users()) {
275  if (auto *I = dyn_cast<LoadInst>(CurUser)) {
276  Loads.push_back({I, Ctx.Offset});
277  continue;
278  }
279 
280  if (auto *I = dyn_cast<BitCastInst>(CurUser)) {
281  Worklist.push({I, Ctx.Offset});
282  continue;
283  }
284 
285  if (auto *I = dyn_cast<GetElementPtrInst>(CurUser)) {
286  APInt OffsetAccumulated =
287  APInt::getZero(DL.getIndexSizeInBits(ADDRESS_SPACE_PARAM));
288 
289  if (!I->accumulateConstantOffset(DL, OffsetAccumulated))
290  continue;
291 
292  uint64_t OffsetLimit = -1;
293  uint64_t Offset = OffsetAccumulated.getLimitedValue(OffsetLimit);
294  assert(Offset != OffsetLimit && "Expect Offset less than UINT64_MAX");
295 
296  Worklist.push({I, Ctx.Offset + Offset});
297  continue;
298  }
299 
300  llvm_unreachable("All users must be one of: load, "
301  "bitcast, getelementptr.");
302  }
303  }
304 
305  for (Load &CurLoad : Loads) {
306  Align NewLoadAlign(greatestCommonDivisor(NewArgAlign, CurLoad.Offset));
307  Align CurLoadAlign(CurLoad.Inst->getAlign());
308  CurLoad.Inst->setAlignment(std::max(NewLoadAlign, CurLoadAlign));
309  }
310 }
311 
312 void NVPTXLowerArgs::handleByValParam(Argument *Arg) {
313  Function *Func = Arg->getParent();
314  Instruction *FirstInst = &(Func->getEntryBlock().front());
315  Type *StructType = Arg->getParamByValType();
316  assert(StructType && "Missing byval type");
317 
318  auto IsALoadChain = [&](Value *Start) {
319  SmallVector<Value *, 16> ValuesToCheck = {Start};
320  auto IsALoadChainInstr = [](Value *V) -> bool {
321  if (isa<GetElementPtrInst>(V) || isa<BitCastInst>(V) || isa<LoadInst>(V))
322  return true;
323  // ASC to param space are OK, too -- we'll just strip them.
324  if (auto *ASC = dyn_cast<AddrSpaceCastInst>(V)) {
325  if (ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM)
326  return true;
327  }
328  return false;
329  };
330 
331  while (!ValuesToCheck.empty()) {
332  Value *V = ValuesToCheck.pop_back_val();
333  if (!IsALoadChainInstr(V)) {
334  LLVM_DEBUG(dbgs() << "Need a copy of " << *Arg << " because of " << *V
335  << "\n");
336  (void)Arg;
337  return false;
338  }
339  if (!isa<LoadInst>(V))
340  llvm::append_range(ValuesToCheck, V->users());
341  }
342  return true;
343  };
344 
345  if (llvm::all_of(Arg->users(), IsALoadChain)) {
346  // Convert all loads and intermediate operations to use parameter AS and
347  // skip creation of a local copy of the argument.
348  SmallVector<User *, 16> UsersToUpdate(Arg->users());
349  Value *ArgInParamAS = new AddrSpaceCastInst(
351  FirstInst);
352  for (Value *V : UsersToUpdate)
353  convertToParamAS(V, ArgInParamAS);
354  LLVM_DEBUG(dbgs() << "No need to copy " << *Arg << "\n");
355 
356  // Further optimizations require target lowering info.
357  if (!TM)
358  return;
359 
360  const auto *TLI =
361  cast<NVPTXTargetLowering>(TM->getSubtargetImpl()->getTargetLowering());
362 
363  adjustByValArgAlignment(Arg, ArgInParamAS, TLI);
364 
365  return;
366  }
367 
368  // Otherwise we have to create a temporary copy.
369  const DataLayout &DL = Func->getParent()->getDataLayout();
370  unsigned AS = DL.getAllocaAddrSpace();
371  AllocaInst *AllocA = new AllocaInst(StructType, AS, Arg->getName(), FirstInst);
372  // Set the alignment to alignment of the byval parameter. This is because,
373  // later load/stores assume that alignment, and we are going to replace
374  // the use of the byval parameter with this alloca instruction.
375  AllocA->setAlignment(Func->getParamAlign(Arg->getArgNo())
376  .value_or(DL.getPrefTypeAlign(StructType)));
377  Arg->replaceAllUsesWith(AllocA);
378 
379  Value *ArgInParam = new AddrSpaceCastInst(
381  FirstInst);
382  // Be sure to propagate alignment to this load; LLVM doesn't know that NVPTX
383  // addrspacecast preserves alignment. Since params are constant, this load is
384  // definitely not volatile.
385  LoadInst *LI =
386  new LoadInst(StructType, ArgInParam, Arg->getName(),
387  /*isVolatile=*/false, AllocA->getAlign(), FirstInst);
388  new StoreInst(LI, AllocA, FirstInst);
389 }
390 
391 void NVPTXLowerArgs::markPointerAsGlobal(Value *Ptr) {
393  return;
394 
395  // Deciding where to emit the addrspacecast pair.
396  BasicBlock::iterator InsertPt;
397  if (Argument *Arg = dyn_cast<Argument>(Ptr)) {
398  // Insert at the functon entry if Ptr is an argument.
399  InsertPt = Arg->getParent()->getEntryBlock().begin();
400  } else {
401  // Insert right after Ptr if Ptr is an instruction.
402  InsertPt = ++cast<Instruction>(Ptr)->getIterator();
403  assert(InsertPt != InsertPt->getParent()->end() &&
404  "We don't call this function with Ptr being a terminator.");
405  }
406 
407  Instruction *PtrInGlobal = new AddrSpaceCastInst(
408  Ptr,
409  PointerType::getWithSamePointeeType(cast<PointerType>(Ptr->getType()),
411  Ptr->getName(), &*InsertPt);
412  Value *PtrInGeneric = new AddrSpaceCastInst(PtrInGlobal, Ptr->getType(),
413  Ptr->getName(), &*InsertPt);
414  // Replace with PtrInGeneric all uses of Ptr except PtrInGlobal.
415  Ptr->replaceAllUsesWith(PtrInGeneric);
416  PtrInGlobal->setOperand(0, Ptr);
417 }
418 
419 // =============================================================================
420 // Main function for this pass.
421 // =============================================================================
422 bool NVPTXLowerArgs::runOnKernelFunction(Function &F) {
423  if (TM && TM->getDrvInterface() == NVPTX::CUDA) {
424  // Mark pointers in byval structs as global.
425  for (auto &B : F) {
426  for (auto &I : B) {
427  if (LoadInst *LI = dyn_cast<LoadInst>(&I)) {
428  if (LI->getType()->isPointerTy()) {
429  Value *UO = getUnderlyingObject(LI->getPointerOperand());
430  if (Argument *Arg = dyn_cast<Argument>(UO)) {
431  if (Arg->hasByValAttr()) {
432  // LI is a load from a pointer within a byval kernel parameter.
433  markPointerAsGlobal(LI);
434  }
435  }
436  }
437  }
438  }
439  }
440  }
441 
442  LLVM_DEBUG(dbgs() << "Lowering kernel args of " << F.getName() << "\n");
443  for (Argument &Arg : F.args()) {
444  if (Arg.getType()->isPointerTy()) {
445  if (Arg.hasByValAttr())
446  handleByValParam(&Arg);
447  else if (TM && TM->getDrvInterface() == NVPTX::CUDA)
448  markPointerAsGlobal(&Arg);
449  }
450  }
451  return true;
452 }
453 
454 // Device functions only need to copy byval args into local memory.
455 bool NVPTXLowerArgs::runOnDeviceFunction(Function &F) {
456  LLVM_DEBUG(dbgs() << "Lowering function args of " << F.getName() << "\n");
457  for (Argument &Arg : F.args())
458  if (Arg.getType()->isPointerTy() && Arg.hasByValAttr())
459  handleByValParam(&Arg);
460  return true;
461 }
462 
464  return isKernelFunction(F) ? runOnKernelFunction(F) : runOnDeviceFunction(F);
465 }
466 
467 FunctionPass *
469  return new NVPTXLowerArgs(TM);
470 }
llvm::Argument
This class represents an incoming formal argument to a Function.
Definition: Argument.h:28
llvm::NVPTXTargetLowering::getFunctionParamOptimizedAlign
Align getFunctionParamOptimizedAlign(const Function *F, Type *ArgTy, const DataLayout &DL) const
getFunctionParamOptimizedAlign - since function arguments are passed via .param space,...
Definition: NVPTXISelLowering.cpp:4265
llvm
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:17
llvm::DataLayout
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:113
llvm::BasicBlock::iterator
InstListType::iterator iterator
Instruction iterators...
Definition: BasicBlock.h:87
llvm::AllocaInst::getAlign
Align getAlign() const
Return the alignment of the memory that is being allocated by the instruction.
Definition: Instructions.h:121
llvm::Function
Definition: Function.h:60
Pass.h
llvm::greatestCommonDivisor
T greatestCommonDivisor(T A, T B)
Return the greatest common divisor of the values using Euclid's algorithm.
Definition: MathExtras.h:648
llvm::PointerType::get
static PointerType * get(Type *ElementType, unsigned AddressSpace)
This constructs a pointer to an object of the specified type in a numbered address space.
Definition: Type.cpp:727
llvm::Attribute::get
static Attribute get(LLVMContext &Context, AttrKind Kind, uint64_t Val=0)
Return a uniquified Attribute object.
Definition: Attributes.cpp:90
llvm::SmallVector
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1185
llvm::Type::getPointerAddressSpace
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
Definition: DerivedTypes.h:729
llvm::CastInst::Create
static CastInst * Create(Instruction::CastOps, Value *S, Type *Ty, const Twine &Name="", Instruction *InsertBefore=nullptr)
Provides a way to construct any of the CastInst subclasses using an opcode instead of the subclass's ...
Definition: Instructions.cpp:3182
llvm::ADDRESS_SPACE_PARAM
@ ADDRESS_SPACE_PARAM
Definition: NVPTXBaseInfo.h:29
InstructionsToDelete
SmallVector< Instruction * > InstructionsToDelete
Definition: NVPTXLowerArgs.cpp:172
ValueTracking.h
llvm::Type
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
Module.h
llvm::reverse
auto reverse(ContainerTy &&C, std::enable_if_t< has_rbegin< ContainerTy >::value > *=nullptr)
Definition: STLExtras.h:380
IP::NewParam
Value * NewParam
Definition: NVPTXLowerArgs.cpp:169
llvm::max
Expected< ExpressionValue > max(const ExpressionValue &Lhs, const ExpressionValue &Rhs)
Definition: FileCheck.cpp:337
llvm::createNVPTXLowerArgsPass
FunctionPass * createNVPTXLowerArgsPass(const NVPTXTargetMachine *TM)
Definition: NVPTXLowerArgs.cpp:468
llvm::SmallVectorImpl::pop_back_val
LLVM_NODISCARD T pop_back_val()
Definition: SmallVector.h:654
llvm::APInt::getZero
static APInt getZero(unsigned numBits)
Get the '0' value for the specified bit-width.
Definition: APInt.h:177
llvm::NVPTXTargetMachine
NVPTXTargetMachine.
Definition: NVPTXTargetMachine.h:25
LLVM_DEBUG
#define LLVM_DEBUG(X)
Definition: Debug.h:101
F
#define F(x, y, z)
Definition: MD5.cpp:55
NVPTX.h
CloneInstInParamAS
auto CloneInstInParamAS
Definition: NVPTXLowerArgs.cpp:174
NVPTXUtilities.h
assert
assert(I &&"OldUser must be an instruction")
llvm::dbgs
raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition: Debug.cpp:163
Arg
amdgpu Simplify well known AMD library false FunctionCallee Value * Arg
Definition: AMDGPULibCalls.cpp:186
llvm::SPII::Load
@ Load
Definition: SparcInstrInfo.h:32
llvm::all_of
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:1617
llvm::AddrSpaceCastInst
This class represents a conversion between pointers from one address space to another.
Definition: Instructions.h:5252
llvm::User
Definition: User.h:44
llvm::ADDRESS_SPACE_GLOBAL
@ ADDRESS_SPACE_GLOBAL
Definition: NVPTXBaseInfo.h:23
Param
Value * Param
Definition: NVPTXLowerArgs.cpp:164
llvm::APInt::getLimitedValue
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:456
IP
Definition: NVPTXLowerArgs.cpp:167
B
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
llvm::Instruction
Definition: Instruction.h:42
IP::OldInstruction
Instruction * OldInstruction
Definition: NVPTXLowerArgs.cpp:168
llvm::PassRegistry
PassRegistry - This class manages the registration and intitialization of the pass subsystem as appli...
Definition: PassRegistry.h:38
llvm::getUnderlyingObject
const Value * getUnderlyingObject(const Value *V, unsigned MaxLookup=6)
This method strips off any GEP address adjustments and pointer casts from the specified value,...
Definition: ValueTracking.cpp:4455
llvm::Align
This struct is a compact representation of a valid (non-zero power of two) alignment.
Definition: Alignment.h:39
llvm::CallingConv::ID
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Definition: CallingConv.h:24
Type.h
llvm::isKernelFunction
bool isKernelFunction(const Function &F)
Definition: NVPTXUtilities.cpp:274
llvm::StoreInst
An instruction for storing to memory.
Definition: Instructions.h:297
uint64_t
NVPTXBaseInfo.h
I
#define I(x, y, z)
Definition: MD5.cpp:58
llvm::APInt
Class for arbitrary precision integers.
Definition: APInt.h:75
llvm::User::setOperand
void setOperand(unsigned i, Value *Val)
Definition: User.h:174
llvm::GetElementPtrInst::Create
static GetElementPtrInst * Create(Type *PointeeType, Value *Ptr, ArrayRef< Value * > IdxList, const Twine &NameStr="", Instruction *InsertBefore=nullptr)
Definition: Instructions.h:942
INITIALIZE_PASS
INITIALIZE_PASS(NVPTXLowerArgs, "nvptx-lower-args", "Lower arguments (NVPTX)", false, false) static void convertToParamAS(Value *OldUser
llvm::StructType
Class to represent struct types.
Definition: DerivedTypes.h:213
llvm::StringRef
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:58
llvm_unreachable
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
Definition: ErrorHandling.h:143
llvm::Value::getType
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:255
llvm::Value::replaceAllUsesWith
void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
Definition: Value.cpp:529
llvm::append_range
void append_range(Container &C, Range &&R)
Wrapper function to append a range to a container.
Definition: STLExtras.h:1823
DL
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
Definition: AArch64SLSHardening.cpp:76
llvm::ifs::IFSSymbolType::Func
@ Func
llvm::Value::getName
StringRef getName() const
Return a constant reference to the value's name.
Definition: Value.cpp:305
llvm::LoadInst
An instruction for reading from memory.
Definition: Instructions.h:173
llvm::NVPTXTargetLowering
Definition: NVPTXISelLowering.h:440
runOnFunction
static bool runOnFunction(Function &F, bool PostInlining)
Definition: EntryExitInstrumenter.cpp:69
llvm::NVPTX::CUDA
@ CUDA
Definition: NVPTX.h:72
llvm::Align::value
uint64_t value() const
This is a hole in the type system and should not be abused.
Definition: Alignment.h:85
adjustByValArgAlignment
static void adjustByValArgAlignment(Argument *Arg, Value *ArgInParamAS, const NVPTXTargetLowering *TLI)
Definition: NVPTXLowerArgs.cpp:234
NVPTXTargetMachine.h
Function.h
llvm::AllocaInst::setAlignment
void setAlignment(Align Align)
Definition: Instructions.h:125
Instructions.h
llvm::PointerType::getWithSamePointeeType
static PointerType * getWithSamePointeeType(PointerType *PT, unsigned AddressSpace)
This constructs a pointer type with the same pointee type as input PointerType (or opaque pointer if ...
Definition: DerivedTypes.h:666
llvm::initializeNVPTXLowerArgsPass
void initializeNVPTXLowerArgsPass(PassRegistry &)
TM
const char LLVMTargetMachineRef TM
Definition: PassBuilderBindings.cpp:47
llvm::FunctionPass
FunctionPass class - This class is used to implement most global optimizations.
Definition: Pass.h:308
GEP
Hexagon Common GEP
Definition: HexagonCommonGEP.cpp:171
ItemsToConvert
SmallVector< IP > ItemsToConvert
Definition: NVPTXLowerArgs.cpp:171
llvm::AllocaInst
an instruction to allocate memory on the stack
Definition: Instructions.h:58
llvm::Value
LLVM Value Representation.
Definition: Value.h:74
llvm::Value::users
iterator_range< user_iterator > users()
Definition: Value.h:421
llvm::Intrinsic::ID
unsigned ID
Definition: TargetTransformInfo.h:38