LLVM  6.0.0svn
NVPTXLowerArgs.cpp
Go to the documentation of this file.
1 //===-- NVPTXLowerArgs.cpp - Lower arguments ------------------------------===//
2 //
3 // The LLVM Compiler Infrastructure
4 //
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
7 //
8 //===----------------------------------------------------------------------===//
9 //
10 //
11 // Arguments to kernel and device functions are passed via param space,
12 // which imposes certain restrictions:
13 // http://docs.nvidia.com/cuda/parallel-thread-execution/#state-spaces
14 //
15 // Kernel parameters are read-only and accessible only via ld.param
16 // instruction, directly or via a pointer. Pointers to kernel
17 // arguments can't be converted to generic address space.
18 //
19 // Device function parameters are directly accessible via
20 // ld.param/st.param, but taking the address of one returns a pointer
21 // to a copy created in local space which *can't* be used with
22 // ld.param/st.param.
23 //
24 // Copying a byval struct into local memory in IR allows us to enforce
25 // the param space restrictions, gives the rest of IR a pointer w/o
26 // param space restrictions, and gives us an opportunity to eliminate
27 // the copy.
28 //
29 // Pointer arguments to kernel functions need more work to be lowered:
30 //
31 // 1. Convert non-byval pointer arguments of CUDA kernels to pointers in the
32 // global address space. This allows later optimizations to emit
33 // ld.global.*/st.global.* for accessing these pointer arguments. For
34 // example,
35 //
36 // define void @foo(float* %input) {
37 // %v = load float, float* %input, align 4
38 // ...
39 // }
40 //
41 // becomes
42 //
43 // define void @foo(float* %input) {
44 // %input2 = addrspacecast float* %input to float addrspace(1)*
45 // %input3 = addrspacecast float addrspace(1)* %input2 to float*
46 // %v = load float, float* %input3, align 4
47 // ...
48 // }
49 //
50 // Later, NVPTXInferAddressSpaces will optimize it to
51 //
52 // define void @foo(float* %input) {
53 // %input2 = addrspacecast float* %input to float addrspace(1)*
54 // %v = load float, float addrspace(1)* %input2, align 4
55 // ...
56 // }
57 //
58 // 2. Convert pointers in a byval kernel parameter to pointers in the global
59 // address space. As #2, it allows NVPTX to emit more ld/st.global. E.g.,
60 //
61 // struct S {
62 // int *x;
63 // int *y;
64 // };
65 // __global__ void foo(S s) {
66 // int *b = s.y;
67 // // use b
68 // }
69 //
70 // "b" points to the global address space. In the IR level,
71 //
72 // define void @foo({i32*, i32*}* byval %input) {
73 // %b_ptr = getelementptr {i32*, i32*}, {i32*, i32*}* %input, i64 0, i32 1
74 // %b = load i32*, i32** %b_ptr
75 // ; use %b
76 // }
77 //
78 // becomes
79 //
80 // define void @foo({i32*, i32*}* byval %input) {
81 // %b_ptr = getelementptr {i32*, i32*}, {i32*, i32*}* %input, i64 0, i32 1
82 // %b = load i32*, i32** %b_ptr
83 // %b_global = addrspacecast i32* %b to i32 addrspace(1)*
84 // %b_generic = addrspacecast i32 addrspace(1)* %b_global to i32*
85 // ; use %b_generic
86 // }
87 //
88 // TODO: merge this pass with NVPTXInferAddressSpaces so that other passes don't
89 // cancel the addrspacecast pair this pass emits.
90 //===----------------------------------------------------------------------===//
91 
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 
102 using namespace llvm;
103 
104 namespace llvm {
106 }
107 
108 namespace {
109 class NVPTXLowerArgs : public FunctionPass {
110  bool runOnFunction(Function &F) override;
111 
112  bool runOnKernelFunction(Function &F);
113  bool runOnDeviceFunction(Function &F);
114 
115  // handle byval parameters
116  void handleByValParam(Argument *Arg);
117  // Knowing Ptr must point to the global address space, this function
118  // addrspacecasts Ptr to global and then back to generic. This allows
119  // NVPTXInferAddressSpaces to fold the global-to-generic cast into
120  // loads/stores that appear later.
121  void markPointerAsGlobal(Value *Ptr);
122 
123 public:
124  static char ID; // Pass identification, replacement for typeid
125  NVPTXLowerArgs(const NVPTXTargetMachine *TM = nullptr)
126  : FunctionPass(ID), TM(TM) {}
127  StringRef getPassName() const override {
128  return "Lower pointer arguments of CUDA kernels";
129  }
130 
131 private:
132  const NVPTXTargetMachine *TM;
133 };
134 } // namespace
135 
136 char NVPTXLowerArgs::ID = 1;
137 
138 INITIALIZE_PASS(NVPTXLowerArgs, "nvptx-lower-args",
139  "Lower arguments (NVPTX)", false, false)
140 
141 // =============================================================================
142 // If the function had a byval struct ptr arg, say foo(%struct.x* byval %d),
143 // then add the following instructions to the first basic block:
144 //
145 // %temp = alloca %struct.x, align 8
146 // %tempd = addrspacecast %struct.x* %d to %struct.x addrspace(101)*
147 // %tv = load %struct.x addrspace(101)* %tempd
148 // store %struct.x %tv, %struct.x* %temp, align 8
149 //
150 // The above code allocates some space in the stack and copies the incoming
151 // struct from param space to local space.
152 // Then replace all occurrences of %d by %temp.
153 // =============================================================================
154 void NVPTXLowerArgs::handleByValParam(Argument *Arg) {
155  Function *Func = Arg->getParent();
156  Instruction *FirstInst = &(Func->getEntryBlock().front());
157  PointerType *PType = dyn_cast<PointerType>(Arg->getType());
158 
159  assert(PType && "Expecting pointer type in handleByValParam");
160 
161  Type *StructType = PType->getElementType();
162  unsigned AS = Func->getParent()->getDataLayout().getAllocaAddrSpace();
163  AllocaInst *AllocA = new AllocaInst(StructType, AS, Arg->getName(), FirstInst);
164  // Set the alignment to alignment of the byval parameter. This is because,
165  // later load/stores assume that alignment, and we are going to replace
166  // the use of the byval parameter with this alloca instruction.
167  AllocA->setAlignment(Func->getParamAlignment(Arg->getArgNo()));
168  Arg->replaceAllUsesWith(AllocA);
169 
170  Value *ArgInParam = new AddrSpaceCastInst(
171  Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(),
172  FirstInst);
173  LoadInst *LI = new LoadInst(ArgInParam, Arg->getName(), FirstInst);
174  new StoreInst(LI, AllocA, FirstInst);
175 }
176 
177 void NVPTXLowerArgs::markPointerAsGlobal(Value *Ptr) {
179  return;
180 
181  // Deciding where to emit the addrspacecast pair.
182  BasicBlock::iterator InsertPt;
183  if (Argument *Arg = dyn_cast<Argument>(Ptr)) {
184  // Insert at the functon entry if Ptr is an argument.
185  InsertPt = Arg->getParent()->getEntryBlock().begin();
186  } else {
187  // Insert right after Ptr if Ptr is an instruction.
188  InsertPt = ++cast<Instruction>(Ptr)->getIterator();
189  assert(InsertPt != InsertPt->getParent()->end() &&
190  "We don't call this function with Ptr being a terminator.");
191  }
192 
193  Instruction *PtrInGlobal = new AddrSpaceCastInst(
196  Ptr->getName(), &*InsertPt);
197  Value *PtrInGeneric = new AddrSpaceCastInst(PtrInGlobal, Ptr->getType(),
198  Ptr->getName(), &*InsertPt);
199  // Replace with PtrInGeneric all uses of Ptr except PtrInGlobal.
200  Ptr->replaceAllUsesWith(PtrInGeneric);
201  PtrInGlobal->setOperand(0, Ptr);
202 }
203 
204 // =============================================================================
205 // Main function for this pass.
206 // =============================================================================
207 bool NVPTXLowerArgs::runOnKernelFunction(Function &F) {
208  if (TM && TM->getDrvInterface() == NVPTX::CUDA) {
209  // Mark pointers in byval structs as global.
210  for (auto &B : F) {
211  for (auto &I : B) {
212  if (LoadInst *LI = dyn_cast<LoadInst>(&I)) {
213  if (LI->getType()->isPointerTy()) {
214  Value *UO = GetUnderlyingObject(LI->getPointerOperand(),
215  F.getParent()->getDataLayout());
216  if (Argument *Arg = dyn_cast<Argument>(UO)) {
217  if (Arg->hasByValAttr()) {
218  // LI is a load from a pointer within a byval kernel parameter.
219  markPointerAsGlobal(LI);
220  }
221  }
222  }
223  }
224  }
225  }
226  }
227 
228  for (Argument &Arg : F.args()) {
229  if (Arg.getType()->isPointerTy()) {
230  if (Arg.hasByValAttr())
231  handleByValParam(&Arg);
232  else if (TM && TM->getDrvInterface() == NVPTX::CUDA)
233  markPointerAsGlobal(&Arg);
234  }
235  }
236  return true;
237 }
238 
239 // Device functions only need to copy byval args into local memory.
240 bool NVPTXLowerArgs::runOnDeviceFunction(Function &F) {
241  for (Argument &Arg : F.args())
242  if (Arg.getType()->isPointerTy() && Arg.hasByValAttr())
243  handleByValParam(&Arg);
244  return true;
245 }
246 
248  return isKernelFunction(F) ? runOnKernelFunction(F) : runOnDeviceFunction(F);
249 }
250 
251 FunctionPass *
253  return new NVPTXLowerArgs(TM);
254 }
void initializeNVPTXLowerArgsPass(PassRegistry &)
This class represents an incoming formal argument to a Function.
Definition: Argument.h:30
Compute iterated dominance frontiers using a linear time algorithm.
Definition: AllocatorList.h:24
void setAlignment(unsigned Align)
unsigned getParamAlignment(unsigned ArgNo) const
Extract the alignment for a call or parameter (0=unknown).
Definition: Function.h:363
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:617
F(f)
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
Definition: DerivedTypes.h:503
An instruction for reading from memory.
Definition: Instructions.h:164
unsigned getAllocaAddrSpace() const
Definition: DataLayout.h:253
This class represents a conversion between pointers from one address space to another.
Type * getPointerElementType() const
Definition: Type.h:373
const DataLayout & getDataLayout() const
Get the data layout for the module&#39;s target platform.
Definition: Module.cpp:361
Class to represent struct types.
Definition: DerivedTypes.h:201
bool isKernelFunction(const Function &F)
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:245
An instruction for storing to memory.
Definition: Instructions.h:306
void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
Definition: Value.cpp:430
Class to represent pointers.
Definition: DerivedTypes.h:467
const BasicBlock & getEntryBlock() const
Definition: Function.h:572
static bool runOnFunction(Function &F, bool PostInlining)
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
The instances of the Type class are immutable: once they are created, they are never changed...
Definition: Type.h:46
bool isPointerTy() const
True if this is an instance of PointerType.
Definition: Type.h:221
const Instruction & front() const
Definition: BasicBlock.h:264
FunctionPass class - This class is used to implement most global optimizations.
Definition: Pass.h:285
const AMDGPUAS & AS
Value * GetUnderlyingObject(Value *V, const DataLayout &DL, unsigned MaxLookup=6)
This method strips off any GEP address adjustments and pointer casts from the specified value...
Iterator for intrusive lists based on ilist_node.
Module.h This file contains the declarations for the Module class.
amdgpu Simplify well known AMD library false Value Value * Arg
NVPTXTargetMachine.
FunctionPass * createNVPTXLowerArgsPass(const NVPTXTargetMachine *TM)
INITIALIZE_PASS(NVPTXLowerArgs, "nvptx-lower-args", "Lower arguments (NVPTX)", false, false) void NVPTXLowerArgs
StringRef getName() const
Return a constant reference to the value&#39;s name.
Definition: Value.cpp:220
#define I(x, y, z)
Definition: MD5.cpp:58
LLVM_NODISCARD std::enable_if<!is_simple_type< Y >::value, typename cast_retty< X, const Y >::ret_type >::type dyn_cast(const Y &Val)
Definition: Casting.h:323
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
Module * getParent()
Get the module that this global value is contained inside of...
Definition: GlobalValue.h:556
LLVM Value Representation.
Definition: Value.h:73
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:49
PassRegistry - This class manages the registration and intitialization of the pass subsystem as appli...
Definition: PassRegistry.h:39
Type * getElementType() const
Definition: DerivedTypes.h:486
iterator_range< arg_iterator > args()
Definition: Function.h:621
an instruction to allocate memory on the stack
Definition: Instructions.h:60