LLVM  14.0.0git
InferAddressSpaces.cpp
Go to the documentation of this file.
1 //===- InferAddressSpace.cpp - --------------------------------------------===//
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 // CUDA C/C++ includes memory space designation as variable type qualifers (such
10 // as __global__ and __shared__). Knowing the space of a memory access allows
11 // CUDA compilers to emit faster PTX loads and stores. For example, a load from
12 // shared memory can be translated to `ld.shared` which is roughly 10% faster
13 // than a generic `ld` on an NVIDIA Tesla K40c.
14 //
15 // Unfortunately, type qualifiers only apply to variable declarations, so CUDA
16 // compilers must infer the memory space of an address expression from
17 // type-qualified variables.
18 //
19 // LLVM IR uses non-zero (so-called) specific address spaces to represent memory
20 // spaces (e.g. addrspace(3) means shared memory). The Clang frontend
21 // places only type-qualified variables in specific address spaces, and then
22 // conservatively `addrspacecast`s each type-qualified variable to addrspace(0)
23 // (so-called the generic address space) for other instructions to use.
24 //
25 // For example, the Clang translates the following CUDA code
26 // __shared__ float a[10];
27 // float v = a[i];
28 // to
29 // %0 = addrspacecast [10 x float] addrspace(3)* @a to [10 x float]*
30 // %1 = gep [10 x float], [10 x float]* %0, i64 0, i64 %i
31 // %v = load float, float* %1 ; emits ld.f32
32 // @a is in addrspace(3) since it's type-qualified, but its use from %1 is
33 // redirected to %0 (the generic version of @a).
34 //
35 // The optimization implemented in this file propagates specific address spaces
36 // from type-qualified variable declarations to its users. For example, it
37 // optimizes the above IR to
38 // %1 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
39 // %v = load float addrspace(3)* %1 ; emits ld.shared.f32
40 // propagating the addrspace(3) from @a to %1. As the result, the NVPTX
41 // codegen is able to emit ld.shared.f32 for %v.
42 //
43 // Address space inference works in two steps. First, it uses a data-flow
44 // analysis to infer as many generic pointers as possible to point to only one
45 // specific address space. In the above example, it can prove that %1 only
46 // points to addrspace(3). This algorithm was published in
47 // CUDA: Compiling and optimizing for a GPU platform
48 // Chakrabarti, Grover, Aarts, Kong, Kudlur, Lin, Marathe, Murphy, Wang
49 // ICCS 2012
50 //
51 // Then, address space inference replaces all refinable generic pointers with
52 // equivalent specific pointers.
53 //
54 // The major challenge of implementing this optimization is handling PHINodes,
55 // which may create loops in the data flow graph. This brings two complications.
56 //
57 // First, the data flow analysis in Step 1 needs to be circular. For example,
58 // %generic.input = addrspacecast float addrspace(3)* %input to float*
59 // loop:
60 // %y = phi [ %generic.input, %y2 ]
61 // %y2 = getelementptr %y, 1
62 // %v = load %y2
63 // br ..., label %loop, ...
64 // proving %y specific requires proving both %generic.input and %y2 specific,
65 // but proving %y2 specific circles back to %y. To address this complication,
66 // the data flow analysis operates on a lattice:
67 // uninitialized > specific address spaces > generic.
68 // All address expressions (our implementation only considers phi, bitcast,
69 // addrspacecast, and getelementptr) start with the uninitialized address space.
70 // The monotone transfer function moves the address space of a pointer down a
71 // lattice path from uninitialized to specific and then to generic. A join
72 // operation of two different specific address spaces pushes the expression down
73 // to the generic address space. The analysis completes once it reaches a fixed
74 // point.
75 //
76 // Second, IR rewriting in Step 2 also needs to be circular. For example,
77 // converting %y to addrspace(3) requires the compiler to know the converted
78 // %y2, but converting %y2 needs the converted %y. To address this complication,
79 // we break these cycles using "undef" placeholders. When converting an
80 // instruction `I` to a new address space, if its operand `Op` is not converted
81 // yet, we let `I` temporarily use `undef` and fix all the uses of undef later.
82 // For instance, our algorithm first converts %y to
83 // %y' = phi float addrspace(3)* [ %input, undef ]
84 // Then, it converts %y2 to
85 // %y2' = getelementptr %y', 1
86 // Finally, it fixes the undef in %y' so that
87 // %y' = phi float addrspace(3)* [ %input, %y2' ]
88 //
89 //===----------------------------------------------------------------------===//
90 
92 #include "llvm/ADT/ArrayRef.h"
93 #include "llvm/ADT/DenseMap.h"
94 #include "llvm/ADT/DenseSet.h"
95 #include "llvm/ADT/None.h"
96 #include "llvm/ADT/Optional.h"
97 #include "llvm/ADT/SetVector.h"
98 #include "llvm/ADT/SmallVector.h"
100 #include "llvm/IR/BasicBlock.h"
101 #include "llvm/IR/Constant.h"
102 #include "llvm/IR/Constants.h"
103 #include "llvm/IR/Function.h"
104 #include "llvm/IR/IRBuilder.h"
105 #include "llvm/IR/InstIterator.h"
106 #include "llvm/IR/Instruction.h"
107 #include "llvm/IR/Instructions.h"
108 #include "llvm/IR/IntrinsicInst.h"
109 #include "llvm/IR/Intrinsics.h"
110 #include "llvm/IR/LLVMContext.h"
111 #include "llvm/IR/Operator.h"
112 #include "llvm/IR/PassManager.h"
113 #include "llvm/IR/Type.h"
114 #include "llvm/IR/Use.h"
115 #include "llvm/IR/User.h"
116 #include "llvm/IR/Value.h"
117 #include "llvm/IR/ValueHandle.h"
118 #include "llvm/Pass.h"
119 #include "llvm/Support/Casting.h"
121 #include "llvm/Support/Compiler.h"
122 #include "llvm/Support/Debug.h"
125 #include "llvm/Transforms/Scalar.h"
128 #include <cassert>
129 #include <iterator>
130 #include <limits>
131 #include <utility>
132 #include <vector>
133 
134 #define DEBUG_TYPE "infer-address-spaces"
135 
136 using namespace llvm;
137 
139  "assume-default-is-flat-addrspace", cl::init(false), cl::ReallyHidden,
140  cl::desc("The default address space is assumed as the flat address space. "
141  "This is mainly for test purpose."));
142 
143 static const unsigned UninitializedAddressSpace =
145 
146 namespace {
147 
148 using ValueToAddrSpaceMapTy = DenseMap<const Value *, unsigned>;
149 using PostorderStackTy = llvm::SmallVector<PointerIntPair<Value *, 1, bool>, 4>;
150 
151 class InferAddressSpaces : public FunctionPass {
152  unsigned FlatAddrSpace = 0;
153 
154 public:
155  static char ID;
156 
157  InferAddressSpaces() :
158  FunctionPass(ID), FlatAddrSpace(UninitializedAddressSpace) {}
159  InferAddressSpaces(unsigned AS) : FunctionPass(ID), FlatAddrSpace(AS) {}
160 
161  void getAnalysisUsage(AnalysisUsage &AU) const override {
162  AU.setPreservesCFG();
164  }
165 
166  bool runOnFunction(Function &F) override;
167 };
168 
169 class InferAddressSpacesImpl {
170  const TargetTransformInfo *TTI = nullptr;
171  const DataLayout *DL = nullptr;
172 
173  /// Target specific address space which uses of should be replaced if
174  /// possible.
175  unsigned FlatAddrSpace = 0;
176 
177  // Returns the new address space of V if updated; otherwise, returns None.
179  updateAddressSpace(const Value &V,
180  const ValueToAddrSpaceMapTy &InferredAddrSpace) const;
181 
182  // Tries to infer the specific address space of each address expression in
183  // Postorder.
184  void inferAddressSpaces(ArrayRef<WeakTrackingVH> Postorder,
185  ValueToAddrSpaceMapTy *InferredAddrSpace) const;
186 
187  bool isSafeToCastConstAddrSpace(Constant *C, unsigned NewAS) const;
188 
189  Value *cloneInstructionWithNewAddressSpace(
190  Instruction *I, unsigned NewAddrSpace,
191  const ValueToValueMapTy &ValueWithNewAddrSpace,
192  SmallVectorImpl<const Use *> *UndefUsesToFix) const;
193 
194  // Changes the flat address expressions in function F to point to specific
195  // address spaces if InferredAddrSpace says so. Postorder is the postorder of
196  // all flat expressions in the use-def graph of function F.
197  bool rewriteWithNewAddressSpaces(
199  const ValueToAddrSpaceMapTy &InferredAddrSpace, Function *F) const;
200 
201  void appendsFlatAddressExpressionToPostorderStack(
202  Value *V, PostorderStackTy &PostorderStack,
203  DenseSet<Value *> &Visited) const;
204 
205  bool rewriteIntrinsicOperands(IntrinsicInst *II,
206  Value *OldV, Value *NewV) const;
207  void collectRewritableIntrinsicOperands(IntrinsicInst *II,
208  PostorderStackTy &PostorderStack,
209  DenseSet<Value *> &Visited) const;
210 
211  std::vector<WeakTrackingVH> collectFlatAddressExpressions(Function &F) const;
212 
213  Value *cloneValueWithNewAddressSpace(
214  Value *V, unsigned NewAddrSpace,
215  const ValueToValueMapTy &ValueWithNewAddrSpace,
216  SmallVectorImpl<const Use *> *UndefUsesToFix) const;
217  unsigned joinAddressSpaces(unsigned AS1, unsigned AS2) const;
218 
219 public:
220  InferAddressSpacesImpl(const TargetTransformInfo *TTI, unsigned FlatAddrSpace)
221  : TTI(TTI), FlatAddrSpace(FlatAddrSpace) {}
222  bool run(Function &F);
223 };
224 
225 } // end anonymous namespace
226 
227 char InferAddressSpaces::ID = 0;
228 
229 namespace llvm {
230 
232 
233 } // end namespace llvm
234 
235 INITIALIZE_PASS(InferAddressSpaces, DEBUG_TYPE, "Infer address spaces",
236  false, false)
237 
238 // Check whether that's no-op pointer bicast using a pair of
239 // `ptrtoint`/`inttoptr` due to the missing no-op pointer bitcast over
240 // different address spaces.
241 static bool isNoopPtrIntCastPair(const Operator *I2P, const DataLayout &DL,
243  assert(I2P->getOpcode() == Instruction::IntToPtr);
244  auto *P2I = dyn_cast<Operator>(I2P->getOperand(0));
245  if (!P2I || P2I->getOpcode() != Instruction::PtrToInt)
246  return false;
247  // Check it's really safe to treat that pair of `ptrtoint`/`inttoptr` as a
248  // no-op cast. Besides checking both of them are no-op casts, as the
249  // reinterpreted pointer may be used in other pointer arithmetic, we also
250  // need to double-check that through the target-specific hook. That ensures
251  // the underlying target also agrees that's a no-op address space cast and
252  // pointer bits are preserved.
253  // The current IR spec doesn't have clear rules on address space casts,
254  // especially a clear definition for pointer bits in non-default address
255  // spaces. It would be undefined if that pointer is dereferenced after an
256  // invalid reinterpret cast. Also, due to the unclearness for the meaning of
257  // bits in non-default address spaces in the current spec, the pointer
258  // arithmetic may also be undefined after invalid pointer reinterpret cast.
259  // However, as we confirm through the target hooks that it's a no-op
260  // addrspacecast, it doesn't matter since the bits should be the same.
261  return CastInst::isNoopCast(Instruction::CastOps(I2P->getOpcode()),
262  I2P->getOperand(0)->getType(), I2P->getType(),
263  DL) &&
265  P2I->getOperand(0)->getType(), P2I->getType(),
266  DL) &&
268  P2I->getOperand(0)->getType()->getPointerAddressSpace(),
269  I2P->getType()->getPointerAddressSpace());
270 }
271 
272 // Returns true if V is an address expression.
273 // TODO: Currently, we consider only phi, bitcast, addrspacecast, and
274 // getelementptr operators.
275 static bool isAddressExpression(const Value &V, const DataLayout &DL,
276  const TargetTransformInfo *TTI) {
277  const Operator *Op = dyn_cast<Operator>(&V);
278  if (!Op)
279  return false;
280 
281  switch (Op->getOpcode()) {
282  case Instruction::PHI:
283  assert(Op->getType()->isPointerTy());
284  return true;
285  case Instruction::BitCast:
286  case Instruction::AddrSpaceCast:
287  case Instruction::GetElementPtr:
288  return true;
289  case Instruction::Select:
290  return Op->getType()->isPointerTy();
291  case Instruction::Call: {
292  const IntrinsicInst *II = dyn_cast<IntrinsicInst>(&V);
293  return II && II->getIntrinsicID() == Intrinsic::ptrmask;
294  }
295  case Instruction::IntToPtr:
296  return isNoopPtrIntCastPair(Op, DL, TTI);
297  default:
298  // That value is an address expression if it has an assumed address space.
300  }
301 }
302 
303 // Returns the pointer operands of V.
304 //
305 // Precondition: V is an address expression.
308  const TargetTransformInfo *TTI) {
309  const Operator &Op = cast<Operator>(V);
310  switch (Op.getOpcode()) {
311  case Instruction::PHI: {
312  auto IncomingValues = cast<PHINode>(Op).incoming_values();
313  return SmallVector<Value *, 2>(IncomingValues.begin(),
314  IncomingValues.end());
315  }
316  case Instruction::BitCast:
317  case Instruction::AddrSpaceCast:
318  case Instruction::GetElementPtr:
319  return {Op.getOperand(0)};
320  case Instruction::Select:
321  return {Op.getOperand(1), Op.getOperand(2)};
322  case Instruction::Call: {
323  const IntrinsicInst &II = cast<IntrinsicInst>(Op);
324  assert(II.getIntrinsicID() == Intrinsic::ptrmask &&
325  "unexpected intrinsic call");
326  return {II.getArgOperand(0)};
327  }
328  case Instruction::IntToPtr: {
329  assert(isNoopPtrIntCastPair(&Op, DL, TTI));
330  auto *P2I = cast<Operator>(Op.getOperand(0));
331  return {P2I->getOperand(0)};
332  }
333  default:
334  llvm_unreachable("Unexpected instruction type.");
335  }
336 }
337 
338 bool InferAddressSpacesImpl::rewriteIntrinsicOperands(IntrinsicInst *II,
339  Value *OldV,
340  Value *NewV) const {
341  Module *M = II->getParent()->getParent()->getParent();
342 
343  switch (II->getIntrinsicID()) {
344  case Intrinsic::objectsize: {
345  Type *DestTy = II->getType();
346  Type *SrcTy = NewV->getType();
347  Function *NewDecl =
348  Intrinsic::getDeclaration(M, II->getIntrinsicID(), {DestTy, SrcTy});
349  II->setArgOperand(0, NewV);
350  II->setCalledFunction(NewDecl);
351  return true;
352  }
353  case Intrinsic::ptrmask:
354  // This is handled as an address expression, not as a use memory operation.
355  return false;
356  default: {
357  Value *Rewrite = TTI->rewriteIntrinsicWithAddressSpace(II, OldV, NewV);
358  if (!Rewrite)
359  return false;
360  if (Rewrite != II)
361  II->replaceAllUsesWith(Rewrite);
362  return true;
363  }
364  }
365 }
366 
367 void InferAddressSpacesImpl::collectRewritableIntrinsicOperands(
368  IntrinsicInst *II, PostorderStackTy &PostorderStack,
369  DenseSet<Value *> &Visited) const {
370  auto IID = II->getIntrinsicID();
371  switch (IID) {
372  case Intrinsic::ptrmask:
373  case Intrinsic::objectsize:
374  appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(0),
375  PostorderStack, Visited);
376  break;
377  default:
378  SmallVector<int, 2> OpIndexes;
379  if (TTI->collectFlatAddressOperands(OpIndexes, IID)) {
380  for (int Idx : OpIndexes) {
381  appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(Idx),
382  PostorderStack, Visited);
383  }
384  }
385  break;
386  }
387 }
388 
389 // Returns all flat address expressions in function F. The elements are
390 // If V is an unvisited flat address expression, appends V to PostorderStack
391 // and marks it as visited.
392 void InferAddressSpacesImpl::appendsFlatAddressExpressionToPostorderStack(
393  Value *V, PostorderStackTy &PostorderStack,
394  DenseSet<Value *> &Visited) const {
395  assert(V->getType()->isPointerTy());
396 
397  // Generic addressing expressions may be hidden in nested constant
398  // expressions.
399  if (ConstantExpr *CE = dyn_cast<ConstantExpr>(V)) {
400  // TODO: Look in non-address parts, like icmp operands.
401  if (isAddressExpression(*CE, *DL, TTI) && Visited.insert(CE).second)
402  PostorderStack.emplace_back(CE, false);
403 
404  return;
405  }
406 
407  if (V->getType()->getPointerAddressSpace() == FlatAddrSpace &&
408  isAddressExpression(*V, *DL, TTI)) {
409  if (Visited.insert(V).second) {
410  PostorderStack.emplace_back(V, false);
411 
412  Operator *Op = cast<Operator>(V);
413  for (unsigned I = 0, E = Op->getNumOperands(); I != E; ++I) {
414  if (ConstantExpr *CE = dyn_cast<ConstantExpr>(Op->getOperand(I))) {
415  if (isAddressExpression(*CE, *DL, TTI) && Visited.insert(CE).second)
416  PostorderStack.emplace_back(CE, false);
417  }
418  }
419  }
420  }
421 }
422 
423 // Returns all flat address expressions in function F. The elements are ordered
424 // ordered in postorder.
425 std::vector<WeakTrackingVH>
426 InferAddressSpacesImpl::collectFlatAddressExpressions(Function &F) const {
427  // This function implements a non-recursive postorder traversal of a partial
428  // use-def graph of function F.
429  PostorderStackTy PostorderStack;
430  // The set of visited expressions.
431  DenseSet<Value *> Visited;
432 
433  auto PushPtrOperand = [&](Value *Ptr) {
434  appendsFlatAddressExpressionToPostorderStack(Ptr, PostorderStack,
435  Visited);
436  };
437 
438  // Look at operations that may be interesting accelerate by moving to a known
439  // address space. We aim at generating after loads and stores, but pure
440  // addressing calculations may also be faster.
441  for (Instruction &I : instructions(F)) {
442  if (auto *GEP = dyn_cast<GetElementPtrInst>(&I)) {
443  if (!GEP->getType()->isVectorTy())
444  PushPtrOperand(GEP->getPointerOperand());
445  } else if (auto *LI = dyn_cast<LoadInst>(&I))
446  PushPtrOperand(LI->getPointerOperand());
447  else if (auto *SI = dyn_cast<StoreInst>(&I))
448  PushPtrOperand(SI->getPointerOperand());
449  else if (auto *RMW = dyn_cast<AtomicRMWInst>(&I))
450  PushPtrOperand(RMW->getPointerOperand());
451  else if (auto *CmpX = dyn_cast<AtomicCmpXchgInst>(&I))
452  PushPtrOperand(CmpX->getPointerOperand());
453  else if (auto *MI = dyn_cast<MemIntrinsic>(&I)) {
454  // For memset/memcpy/memmove, any pointer operand can be replaced.
455  PushPtrOperand(MI->getRawDest());
456 
457  // Handle 2nd operand for memcpy/memmove.
458  if (auto *MTI = dyn_cast<MemTransferInst>(MI))
459  PushPtrOperand(MTI->getRawSource());
460  } else if (auto *II = dyn_cast<IntrinsicInst>(&I))
461  collectRewritableIntrinsicOperands(II, PostorderStack, Visited);
462  else if (ICmpInst *Cmp = dyn_cast<ICmpInst>(&I)) {
463  // FIXME: Handle vectors of pointers
464  if (Cmp->getOperand(0)->getType()->isPointerTy()) {
465  PushPtrOperand(Cmp->getOperand(0));
466  PushPtrOperand(Cmp->getOperand(1));
467  }
468  } else if (auto *ASC = dyn_cast<AddrSpaceCastInst>(&I)) {
469  if (!ASC->getType()->isVectorTy())
470  PushPtrOperand(ASC->getPointerOperand());
471  } else if (auto *I2P = dyn_cast<IntToPtrInst>(&I)) {
472  if (isNoopPtrIntCastPair(cast<Operator>(I2P), *DL, TTI))
473  PushPtrOperand(
474  cast<Operator>(I2P->getOperand(0))->getOperand(0));
475  }
476  }
477 
478  std::vector<WeakTrackingVH> Postorder; // The resultant postorder.
479  while (!PostorderStack.empty()) {
480  Value *TopVal = PostorderStack.back().getPointer();
481  // If the operands of the expression on the top are already explored,
482  // adds that expression to the resultant postorder.
483  if (PostorderStack.back().getInt()) {
484  if (TopVal->getType()->getPointerAddressSpace() == FlatAddrSpace)
485  Postorder.push_back(TopVal);
486  PostorderStack.pop_back();
487  continue;
488  }
489  // Otherwise, adds its operands to the stack and explores them.
490  PostorderStack.back().setInt(true);
491  // Skip values with an assumed address space.
493  for (Value *PtrOperand : getPointerOperands(*TopVal, *DL, TTI)) {
494  appendsFlatAddressExpressionToPostorderStack(PtrOperand, PostorderStack,
495  Visited);
496  }
497  }
498  }
499  return Postorder;
500 }
501 
502 // A helper function for cloneInstructionWithNewAddressSpace. Returns the clone
503 // of OperandUse.get() in the new address space. If the clone is not ready yet,
504 // returns an undef in the new address space as a placeholder.
506  const Use &OperandUse, unsigned NewAddrSpace,
507  const ValueToValueMapTy &ValueWithNewAddrSpace,
508  SmallVectorImpl<const Use *> *UndefUsesToFix) {
509  Value *Operand = OperandUse.get();
510 
512  cast<PointerType>(Operand->getType()), NewAddrSpace);
513 
514  if (Constant *C = dyn_cast<Constant>(Operand))
515  return ConstantExpr::getAddrSpaceCast(C, NewPtrTy);
516 
517  if (Value *NewOperand = ValueWithNewAddrSpace.lookup(Operand))
518  return NewOperand;
519 
520  UndefUsesToFix->push_back(&OperandUse);
521  return UndefValue::get(NewPtrTy);
522 }
523 
524 // Returns a clone of `I` with its operands converted to those specified in
525 // ValueWithNewAddrSpace. Due to potential cycles in the data flow graph, an
526 // operand whose address space needs to be modified might not exist in
527 // ValueWithNewAddrSpace. In that case, uses undef as a placeholder operand and
528 // adds that operand use to UndefUsesToFix so that caller can fix them later.
529 //
530 // Note that we do not necessarily clone `I`, e.g., if it is an addrspacecast
531 // from a pointer whose type already matches. Therefore, this function returns a
532 // Value* instead of an Instruction*.
533 //
534 // This may also return nullptr in the case the instruction could not be
535 // rewritten.
536 Value *InferAddressSpacesImpl::cloneInstructionWithNewAddressSpace(
537  Instruction *I, unsigned NewAddrSpace,
538  const ValueToValueMapTy &ValueWithNewAddrSpace,
539  SmallVectorImpl<const Use *> *UndefUsesToFix) const {
541  cast<PointerType>(I->getType()), NewAddrSpace);
542 
543  if (I->getOpcode() == Instruction::AddrSpaceCast) {
544  Value *Src = I->getOperand(0);
545  // Because `I` is flat, the source address space must be specific.
546  // Therefore, the inferred address space must be the source space, according
547  // to our algorithm.
548  assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
549  if (Src->getType() != NewPtrType)
550  return new BitCastInst(Src, NewPtrType);
551  return Src;
552  }
553 
554  if (IntrinsicInst *II = dyn_cast<IntrinsicInst>(I)) {
555  // Technically the intrinsic ID is a pointer typed argument, so specially
556  // handle calls early.
557  assert(II->getIntrinsicID() == Intrinsic::ptrmask);
559  II->getArgOperandUse(0), NewAddrSpace, ValueWithNewAddrSpace,
560  UndefUsesToFix);
561  Value *Rewrite =
563  if (Rewrite) {
564  assert(Rewrite != II && "cannot modify this pointer operation in place");
565  return Rewrite;
566  }
567 
568  return nullptr;
569  }
570 
571  unsigned AS = TTI->getAssumedAddrSpace(I);
572  if (AS != UninitializedAddressSpace) {
573  // For the assumed address space, insert an `addrspacecast` to make that
574  // explicit.
576  cast<PointerType>(I->getType()), AS);
577  auto *NewI = new AddrSpaceCastInst(I, NewPtrTy);
578  NewI->insertAfter(I);
579  return NewI;
580  }
581 
582  // Computes the converted pointer operands.
583  SmallVector<Value *, 4> NewPointerOperands;
584  for (const Use &OperandUse : I->operands()) {
585  if (!OperandUse.get()->getType()->isPointerTy())
586  NewPointerOperands.push_back(nullptr);
587  else
588  NewPointerOperands.push_back(operandWithNewAddressSpaceOrCreateUndef(
589  OperandUse, NewAddrSpace, ValueWithNewAddrSpace, UndefUsesToFix));
590  }
591 
592  switch (I->getOpcode()) {
593  case Instruction::BitCast:
594  return new BitCastInst(NewPointerOperands[0], NewPtrType);
595  case Instruction::PHI: {
596  assert(I->getType()->isPointerTy());
597  PHINode *PHI = cast<PHINode>(I);
598  PHINode *NewPHI = PHINode::Create(NewPtrType, PHI->getNumIncomingValues());
599  for (unsigned Index = 0; Index < PHI->getNumIncomingValues(); ++Index) {
600  unsigned OperandNo = PHINode::getOperandNumForIncomingValue(Index);
601  NewPHI->addIncoming(NewPointerOperands[OperandNo],
602  PHI->getIncomingBlock(Index));
603  }
604  return NewPHI;
605  }
606  case Instruction::GetElementPtr: {
607  GetElementPtrInst *GEP = cast<GetElementPtrInst>(I);
609  GEP->getSourceElementType(), NewPointerOperands[0],
610  SmallVector<Value *, 4>(GEP->indices()));
611  NewGEP->setIsInBounds(GEP->isInBounds());
612  return NewGEP;
613  }
614  case Instruction::Select:
615  assert(I->getType()->isPointerTy());
616  return SelectInst::Create(I->getOperand(0), NewPointerOperands[1],
617  NewPointerOperands[2], "", nullptr, I);
618  case Instruction::IntToPtr: {
619  assert(isNoopPtrIntCastPair(cast<Operator>(I), *DL, TTI));
620  Value *Src = cast<Operator>(I->getOperand(0))->getOperand(0);
621  assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
622  if (Src->getType() != NewPtrType)
623  return new BitCastInst(Src, NewPtrType);
624  return Src;
625  }
626  default:
627  llvm_unreachable("Unexpected opcode");
628  }
629 }
630 
631 // Similar to cloneInstructionWithNewAddressSpace, returns a clone of the
632 // constant expression `CE` with its operands replaced as specified in
633 // ValueWithNewAddrSpace.
635  ConstantExpr *CE, unsigned NewAddrSpace,
636  const ValueToValueMapTy &ValueWithNewAddrSpace, const DataLayout *DL,
637  const TargetTransformInfo *TTI) {
638  Type *TargetType = CE->getType()->isPointerTy()
640  cast<PointerType>(CE->getType()), NewAddrSpace)
641  : CE->getType();
642 
643  if (CE->getOpcode() == Instruction::AddrSpaceCast) {
644  // Because CE is flat, the source address space must be specific.
645  // Therefore, the inferred address space must be the source space according
646  // to our algorithm.
647  assert(CE->getOperand(0)->getType()->getPointerAddressSpace() ==
648  NewAddrSpace);
649  return ConstantExpr::getBitCast(CE->getOperand(0), TargetType);
650  }
651 
652  if (CE->getOpcode() == Instruction::BitCast) {
653  if (Value *NewOperand = ValueWithNewAddrSpace.lookup(CE->getOperand(0)))
654  return ConstantExpr::getBitCast(cast<Constant>(NewOperand), TargetType);
655  return ConstantExpr::getAddrSpaceCast(CE, TargetType);
656  }
657 
658  if (CE->getOpcode() == Instruction::Select) {
659  Constant *Src0 = CE->getOperand(1);
660  Constant *Src1 = CE->getOperand(2);
661  if (Src0->getType()->getPointerAddressSpace() ==
662  Src1->getType()->getPointerAddressSpace()) {
663 
665  CE->getOperand(0), ConstantExpr::getAddrSpaceCast(Src0, TargetType),
666  ConstantExpr::getAddrSpaceCast(Src1, TargetType));
667  }
668  }
669 
670  if (CE->getOpcode() == Instruction::IntToPtr) {
671  assert(isNoopPtrIntCastPair(cast<Operator>(CE), *DL, TTI));
672  Constant *Src = cast<ConstantExpr>(CE->getOperand(0))->getOperand(0);
673  assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
674  return ConstantExpr::getBitCast(Src, TargetType);
675  }
676 
677  // Computes the operands of the new constant expression.
678  bool IsNew = false;
679  SmallVector<Constant *, 4> NewOperands;
680  for (unsigned Index = 0; Index < CE->getNumOperands(); ++Index) {
681  Constant *Operand = CE->getOperand(Index);
682  // If the address space of `Operand` needs to be modified, the new operand
683  // with the new address space should already be in ValueWithNewAddrSpace
684  // because (1) the constant expressions we consider (i.e. addrspacecast,
685  // bitcast, and getelementptr) do not incur cycles in the data flow graph
686  // and (2) this function is called on constant expressions in postorder.
687  if (Value *NewOperand = ValueWithNewAddrSpace.lookup(Operand)) {
688  IsNew = true;
689  NewOperands.push_back(cast<Constant>(NewOperand));
690  continue;
691  }
692  if (auto CExpr = dyn_cast<ConstantExpr>(Operand))
694  CExpr, NewAddrSpace, ValueWithNewAddrSpace, DL, TTI)) {
695  IsNew = true;
696  NewOperands.push_back(cast<Constant>(NewOperand));
697  continue;
698  }
699  // Otherwise, reuses the old operand.
700  NewOperands.push_back(Operand);
701  }
702 
703  // If !IsNew, we will replace the Value with itself. However, replaced values
704  // are assumed to wrapped in a addrspace cast later so drop it now.
705  if (!IsNew)
706  return nullptr;
707 
708  if (CE->getOpcode() == Instruction::GetElementPtr) {
709  // Needs to specify the source type while constructing a getelementptr
710  // constant expression.
711  return CE->getWithOperands(
712  NewOperands, TargetType, /*OnlyIfReduced=*/false,
713  NewOperands[0]->getType()->getPointerElementType());
714  }
715 
716  return CE->getWithOperands(NewOperands, TargetType);
717 }
718 
719 // Returns a clone of the value `V`, with its operands replaced as specified in
720 // ValueWithNewAddrSpace. This function is called on every flat address
721 // expression whose address space needs to be modified, in postorder.
722 //
723 // See cloneInstructionWithNewAddressSpace for the meaning of UndefUsesToFix.
724 Value *InferAddressSpacesImpl::cloneValueWithNewAddressSpace(
725  Value *V, unsigned NewAddrSpace,
726  const ValueToValueMapTy &ValueWithNewAddrSpace,
727  SmallVectorImpl<const Use *> *UndefUsesToFix) const {
728  // All values in Postorder are flat address expressions.
729  assert(V->getType()->getPointerAddressSpace() == FlatAddrSpace &&
730  isAddressExpression(*V, *DL, TTI));
731 
732  if (Instruction *I = dyn_cast<Instruction>(V)) {
733  Value *NewV = cloneInstructionWithNewAddressSpace(
734  I, NewAddrSpace, ValueWithNewAddrSpace, UndefUsesToFix);
735  if (Instruction *NewI = dyn_cast_or_null<Instruction>(NewV)) {
736  if (NewI->getParent() == nullptr) {
737  NewI->insertBefore(I);
738  NewI->takeName(I);
739  }
740  }
741  return NewV;
742  }
743 
745  cast<ConstantExpr>(V), NewAddrSpace, ValueWithNewAddrSpace, DL, TTI);
746 }
747 
748 // Defines the join operation on the address space lattice (see the file header
749 // comments).
750 unsigned InferAddressSpacesImpl::joinAddressSpaces(unsigned AS1,
751  unsigned AS2) const {
752  if (AS1 == FlatAddrSpace || AS2 == FlatAddrSpace)
753  return FlatAddrSpace;
754 
755  if (AS1 == UninitializedAddressSpace)
756  return AS2;
757  if (AS2 == UninitializedAddressSpace)
758  return AS1;
759 
760  // The join of two different specific address spaces is flat.
761  return (AS1 == AS2) ? AS1 : FlatAddrSpace;
762 }
763 
764 bool InferAddressSpacesImpl::run(Function &F) {
765  DL = &F.getParent()->getDataLayout();
766 
768  FlatAddrSpace = 0;
769 
770  if (FlatAddrSpace == UninitializedAddressSpace) {
771  FlatAddrSpace = TTI->getFlatAddressSpace();
772  if (FlatAddrSpace == UninitializedAddressSpace)
773  return false;
774  }
775 
776  // Collects all flat address expressions in postorder.
777  std::vector<WeakTrackingVH> Postorder = collectFlatAddressExpressions(F);
778 
779  // Runs a data-flow analysis to refine the address spaces of every expression
780  // in Postorder.
781  ValueToAddrSpaceMapTy InferredAddrSpace;
782  inferAddressSpaces(Postorder, &InferredAddrSpace);
783 
784  // Changes the address spaces of the flat address expressions who are inferred
785  // to point to a specific address space.
786  return rewriteWithNewAddressSpaces(*TTI, Postorder, InferredAddrSpace, &F);
787 }
788 
789 // Constants need to be tracked through RAUW to handle cases with nested
790 // constant expressions, so wrap values in WeakTrackingVH.
791 void InferAddressSpacesImpl::inferAddressSpaces(
792  ArrayRef<WeakTrackingVH> Postorder,
793  ValueToAddrSpaceMapTy *InferredAddrSpace) const {
794  SetVector<Value *> Worklist(Postorder.begin(), Postorder.end());
795  // Initially, all expressions are in the uninitialized address space.
796  for (Value *V : Postorder)
797  (*InferredAddrSpace)[V] = UninitializedAddressSpace;
798 
799  while (!Worklist.empty()) {
800  Value *V = Worklist.pop_back_val();
801 
802  // Tries to update the address space of the stack top according to the
803  // address spaces of its operands.
804  LLVM_DEBUG(dbgs() << "Updating the address space of\n " << *V << '\n');
805  Optional<unsigned> NewAS = updateAddressSpace(*V, *InferredAddrSpace);
806  if (!NewAS.hasValue())
807  continue;
808  // If any updates are made, grabs its users to the worklist because
809  // their address spaces can also be possibly updated.
810  LLVM_DEBUG(dbgs() << " to " << NewAS.getValue() << '\n');
811  (*InferredAddrSpace)[V] = NewAS.getValue();
812 
813  for (Value *User : V->users()) {
814  // Skip if User is already in the worklist.
815  if (Worklist.count(User))
816  continue;
817 
818  auto Pos = InferredAddrSpace->find(User);
819  // Our algorithm only updates the address spaces of flat address
820  // expressions, which are those in InferredAddrSpace.
821  if (Pos == InferredAddrSpace->end())
822  continue;
823 
824  // Function updateAddressSpace moves the address space down a lattice
825  // path. Therefore, nothing to do if User is already inferred as flat (the
826  // bottom element in the lattice).
827  if (Pos->second == FlatAddrSpace)
828  continue;
829 
830  Worklist.insert(User);
831  }
832  }
833 }
834 
835 Optional<unsigned> InferAddressSpacesImpl::updateAddressSpace(
836  const Value &V, const ValueToAddrSpaceMapTy &InferredAddrSpace) const {
837  assert(InferredAddrSpace.count(&V));
838 
839  // The new inferred address space equals the join of the address spaces
840  // of all its pointer operands.
841  unsigned NewAS = UninitializedAddressSpace;
842 
843  const Operator &Op = cast<Operator>(V);
844  if (Op.getOpcode() == Instruction::Select) {
845  Value *Src0 = Op.getOperand(1);
846  Value *Src1 = Op.getOperand(2);
847 
848  auto I = InferredAddrSpace.find(Src0);
849  unsigned Src0AS = (I != InferredAddrSpace.end()) ?
850  I->second : Src0->getType()->getPointerAddressSpace();
851 
852  auto J = InferredAddrSpace.find(Src1);
853  unsigned Src1AS = (J != InferredAddrSpace.end()) ?
854  J->second : Src1->getType()->getPointerAddressSpace();
855 
856  auto *C0 = dyn_cast<Constant>(Src0);
857  auto *C1 = dyn_cast<Constant>(Src1);
858 
859  // If one of the inputs is a constant, we may be able to do a constant
860  // addrspacecast of it. Defer inferring the address space until the input
861  // address space is known.
862  if ((C1 && Src0AS == UninitializedAddressSpace) ||
863  (C0 && Src1AS == UninitializedAddressSpace))
864  return None;
865 
866  if (C0 && isSafeToCastConstAddrSpace(C0, Src1AS))
867  NewAS = Src1AS;
868  else if (C1 && isSafeToCastConstAddrSpace(C1, Src0AS))
869  NewAS = Src0AS;
870  else
871  NewAS = joinAddressSpaces(Src0AS, Src1AS);
872  } else {
873  unsigned AS = TTI->getAssumedAddrSpace(&V);
874  if (AS != UninitializedAddressSpace) {
875  // Use the assumed address space directly.
876  NewAS = AS;
877  } else {
878  // Otherwise, infer the address space from its pointer operands.
879  for (Value *PtrOperand : getPointerOperands(V, *DL, TTI)) {
880  auto I = InferredAddrSpace.find(PtrOperand);
881  unsigned OperandAS =
882  I != InferredAddrSpace.end()
883  ? I->second
884  : PtrOperand->getType()->getPointerAddressSpace();
885 
886  // join(flat, *) = flat. So we can break if NewAS is already flat.
887  NewAS = joinAddressSpaces(NewAS, OperandAS);
888  if (NewAS == FlatAddrSpace)
889  break;
890  }
891  }
892  }
893 
894  unsigned OldAS = InferredAddrSpace.lookup(&V);
895  assert(OldAS != FlatAddrSpace);
896  if (OldAS == NewAS)
897  return None;
898  return NewAS;
899 }
900 
901 /// \p returns true if \p U is the pointer operand of a memory instruction with
902 /// a single pointer operand that can have its address space changed by simply
903 /// mutating the use to a new value. If the memory instruction is volatile,
904 /// return true only if the target allows the memory instruction to be volatile
905 /// in the new address space.
907  Use &U, unsigned AddrSpace) {
908  User *Inst = U.getUser();
909  unsigned OpNo = U.getOperandNo();
910  bool VolatileIsAllowed = false;
911  if (auto *I = dyn_cast<Instruction>(Inst))
912  VolatileIsAllowed = TTI.hasVolatileVariant(I, AddrSpace);
913 
914  if (auto *LI = dyn_cast<LoadInst>(Inst))
915  return OpNo == LoadInst::getPointerOperandIndex() &&
916  (VolatileIsAllowed || !LI->isVolatile());
917 
918  if (auto *SI = dyn_cast<StoreInst>(Inst))
919  return OpNo == StoreInst::getPointerOperandIndex() &&
920  (VolatileIsAllowed || !SI->isVolatile());
921 
922  if (auto *RMW = dyn_cast<AtomicRMWInst>(Inst))
923  return OpNo == AtomicRMWInst::getPointerOperandIndex() &&
924  (VolatileIsAllowed || !RMW->isVolatile());
925 
926  if (auto *CmpX = dyn_cast<AtomicCmpXchgInst>(Inst))
928  (VolatileIsAllowed || !CmpX->isVolatile());
929 
930  return false;
931 }
932 
933 /// Update memory intrinsic uses that require more complex processing than
934 /// simple memory instructions. Thse require re-mangling and may have multiple
935 /// pointer operands.
937  Value *NewV) {
938  IRBuilder<> B(MI);
939  MDNode *TBAA = MI->getMetadata(LLVMContext::MD_tbaa);
940  MDNode *ScopeMD = MI->getMetadata(LLVMContext::MD_alias_scope);
941  MDNode *NoAliasMD = MI->getMetadata(LLVMContext::MD_noalias);
942 
943  if (auto *MSI = dyn_cast<MemSetInst>(MI)) {
944  B.CreateMemSet(NewV, MSI->getValue(), MSI->getLength(),
945  MaybeAlign(MSI->getDestAlignment()),
946  false, // isVolatile
947  TBAA, ScopeMD, NoAliasMD);
948  } else if (auto *MTI = dyn_cast<MemTransferInst>(MI)) {
949  Value *Src = MTI->getRawSource();
950  Value *Dest = MTI->getRawDest();
951 
952  // Be careful in case this is a self-to-self copy.
953  if (Src == OldV)
954  Src = NewV;
955 
956  if (Dest == OldV)
957  Dest = NewV;
958 
959  if (isa<MemCpyInlineInst>(MTI)) {
960  MDNode *TBAAStruct = MTI->getMetadata(LLVMContext::MD_tbaa_struct);
961  B.CreateMemCpyInline(Dest, MTI->getDestAlign(), Src,
962  MTI->getSourceAlign(), MTI->getLength(),
963  false, // isVolatile
964  TBAA, TBAAStruct, ScopeMD, NoAliasMD);
965  } else if (isa<MemCpyInst>(MTI)) {
966  MDNode *TBAAStruct = MTI->getMetadata(LLVMContext::MD_tbaa_struct);
967  B.CreateMemCpy(Dest, MTI->getDestAlign(), Src, MTI->getSourceAlign(),
968  MTI->getLength(),
969  false, // isVolatile
970  TBAA, TBAAStruct, ScopeMD, NoAliasMD);
971  } else {
972  assert(isa<MemMoveInst>(MTI));
973  B.CreateMemMove(Dest, MTI->getDestAlign(), Src, MTI->getSourceAlign(),
974  MTI->getLength(),
975  false, // isVolatile
976  TBAA, ScopeMD, NoAliasMD);
977  }
978  } else
979  llvm_unreachable("unhandled MemIntrinsic");
980 
981  MI->eraseFromParent();
982  return true;
983 }
984 
985 // \p returns true if it is OK to change the address space of constant \p C with
986 // a ConstantExpr addrspacecast.
987 bool InferAddressSpacesImpl::isSafeToCastConstAddrSpace(Constant *C,
988  unsigned NewAS) const {
990 
991  unsigned SrcAS = C->getType()->getPointerAddressSpace();
992  if (SrcAS == NewAS || isa<UndefValue>(C))
993  return true;
994 
995  // Prevent illegal casts between different non-flat address spaces.
996  if (SrcAS != FlatAddrSpace && NewAS != FlatAddrSpace)
997  return false;
998 
999  if (isa<ConstantPointerNull>(C))
1000  return true;
1001 
1002  if (auto *Op = dyn_cast<Operator>(C)) {
1003  // If we already have a constant addrspacecast, it should be safe to cast it
1004  // off.
1005  if (Op->getOpcode() == Instruction::AddrSpaceCast)
1006  return isSafeToCastConstAddrSpace(cast<Constant>(Op->getOperand(0)), NewAS);
1007 
1008  if (Op->getOpcode() == Instruction::IntToPtr &&
1009  Op->getType()->getPointerAddressSpace() == FlatAddrSpace)
1010  return true;
1011  }
1012 
1013  return false;
1014 }
1015 
1017  Value::use_iterator End) {
1018  User *CurUser = I->getUser();
1019  ++I;
1020 
1021  while (I != End && I->getUser() == CurUser)
1022  ++I;
1023 
1024  return I;
1025 }
1026 
1027 bool InferAddressSpacesImpl::rewriteWithNewAddressSpaces(
1029  const ValueToAddrSpaceMapTy &InferredAddrSpace, Function *F) const {
1030  // For each address expression to be modified, creates a clone of it with its
1031  // pointer operands converted to the new address space. Since the pointer
1032  // operands are converted, the clone is naturally in the new address space by
1033  // construction.
1034  ValueToValueMapTy ValueWithNewAddrSpace;
1035  SmallVector<const Use *, 32> UndefUsesToFix;
1036  for (Value* V : Postorder) {
1037  unsigned NewAddrSpace = InferredAddrSpace.lookup(V);
1038 
1039  // In some degenerate cases (e.g. invalid IR in unreachable code), we may
1040  // not even infer the value to have its original address space.
1041  if (NewAddrSpace == UninitializedAddressSpace)
1042  continue;
1043 
1044  if (V->getType()->getPointerAddressSpace() != NewAddrSpace) {
1045  Value *New = cloneValueWithNewAddressSpace(
1046  V, NewAddrSpace, ValueWithNewAddrSpace, &UndefUsesToFix);
1047  if (New)
1048  ValueWithNewAddrSpace[V] = New;
1049  }
1050  }
1051 
1052  if (ValueWithNewAddrSpace.empty())
1053  return false;
1054 
1055  // Fixes all the undef uses generated by cloneInstructionWithNewAddressSpace.
1056  for (const Use *UndefUse : UndefUsesToFix) {
1057  User *V = UndefUse->getUser();
1058  User *NewV = cast_or_null<User>(ValueWithNewAddrSpace.lookup(V));
1059  if (!NewV)
1060  continue;
1061 
1062  unsigned OperandNo = UndefUse->getOperandNo();
1063  assert(isa<UndefValue>(NewV->getOperand(OperandNo)));
1064  NewV->setOperand(OperandNo, ValueWithNewAddrSpace.lookup(UndefUse->get()));
1065  }
1066 
1067  SmallVector<Instruction *, 16> DeadInstructions;
1068 
1069  // Replaces the uses of the old address expressions with the new ones.
1070  for (const WeakTrackingVH &WVH : Postorder) {
1071  assert(WVH && "value was unexpectedly deleted");
1072  Value *V = WVH;
1073  Value *NewV = ValueWithNewAddrSpace.lookup(V);
1074  if (NewV == nullptr)
1075  continue;
1076 
1077  LLVM_DEBUG(dbgs() << "Replacing the uses of " << *V << "\n with\n "
1078  << *NewV << '\n');
1079 
1080  if (Constant *C = dyn_cast<Constant>(V)) {
1081  Constant *Replace = ConstantExpr::getAddrSpaceCast(cast<Constant>(NewV),
1082  C->getType());
1083  if (C != Replace) {
1084  LLVM_DEBUG(dbgs() << "Inserting replacement const cast: " << Replace
1085  << ": " << *Replace << '\n');
1086  C->replaceAllUsesWith(Replace);
1087  V = Replace;
1088  }
1089  }
1090 
1091  Value::use_iterator I, E, Next;
1092  for (I = V->use_begin(), E = V->use_end(); I != E; ) {
1093  Use &U = *I;
1094 
1095  // Some users may see the same pointer operand in multiple operands. Skip
1096  // to the next instruction.
1097  I = skipToNextUser(I, E);
1098 
1100  TTI, U, V->getType()->getPointerAddressSpace())) {
1101  // If V is used as the pointer operand of a compatible memory operation,
1102  // sets the pointer operand to NewV. This replacement does not change
1103  // the element type, so the resultant load/store is still valid.
1104  U.set(NewV);
1105  continue;
1106  }
1107 
1108  User *CurUser = U.getUser();
1109  // Skip if the current user is the new value itself.
1110  if (CurUser == NewV)
1111  continue;
1112  // Handle more complex cases like intrinsic that need to be remangled.
1113  if (auto *MI = dyn_cast<MemIntrinsic>(CurUser)) {
1114  if (!MI->isVolatile() && handleMemIntrinsicPtrUse(MI, V, NewV))
1115  continue;
1116  }
1117 
1118  if (auto *II = dyn_cast<IntrinsicInst>(CurUser)) {
1119  if (rewriteIntrinsicOperands(II, V, NewV))
1120  continue;
1121  }
1122 
1123  if (isa<Instruction>(CurUser)) {
1124  if (ICmpInst *Cmp = dyn_cast<ICmpInst>(CurUser)) {
1125  // If we can infer that both pointers are in the same addrspace,
1126  // transform e.g.
1127  // %cmp = icmp eq float* %p, %q
1128  // into
1129  // %cmp = icmp eq float addrspace(3)* %new_p, %new_q
1130 
1131  unsigned NewAS = NewV->getType()->getPointerAddressSpace();
1132  int SrcIdx = U.getOperandNo();
1133  int OtherIdx = (SrcIdx == 0) ? 1 : 0;
1134  Value *OtherSrc = Cmp->getOperand(OtherIdx);
1135 
1136  if (Value *OtherNewV = ValueWithNewAddrSpace.lookup(OtherSrc)) {
1137  if (OtherNewV->getType()->getPointerAddressSpace() == NewAS) {
1138  Cmp->setOperand(OtherIdx, OtherNewV);
1139  Cmp->setOperand(SrcIdx, NewV);
1140  continue;
1141  }
1142  }
1143 
1144  // Even if the type mismatches, we can cast the constant.
1145  if (auto *KOtherSrc = dyn_cast<Constant>(OtherSrc)) {
1146  if (isSafeToCastConstAddrSpace(KOtherSrc, NewAS)) {
1147  Cmp->setOperand(SrcIdx, NewV);
1148  Cmp->setOperand(OtherIdx,
1149  ConstantExpr::getAddrSpaceCast(KOtherSrc, NewV->getType()));
1150  continue;
1151  }
1152  }
1153  }
1154 
1155  if (AddrSpaceCastInst *ASC = dyn_cast<AddrSpaceCastInst>(CurUser)) {
1156  unsigned NewAS = NewV->getType()->getPointerAddressSpace();
1157  if (ASC->getDestAddressSpace() == NewAS) {
1158  if (ASC->getType()->getPointerElementType() !=
1159  NewV->getType()->getPointerElementType()) {
1160  NewV = CastInst::Create(Instruction::BitCast, NewV,
1161  ASC->getType(), "", ASC);
1162  }
1163  ASC->replaceAllUsesWith(NewV);
1164  DeadInstructions.push_back(ASC);
1165  continue;
1166  }
1167  }
1168 
1169  // Otherwise, replaces the use with flat(NewV).
1170  if (Instruction *Inst = dyn_cast<Instruction>(V)) {
1171  // Don't create a copy of the original addrspacecast.
1172  if (U == V && isa<AddrSpaceCastInst>(V))
1173  continue;
1174 
1175  BasicBlock::iterator InsertPos = std::next(Inst->getIterator());
1176  while (isa<PHINode>(InsertPos))
1177  ++InsertPos;
1178  U.set(new AddrSpaceCastInst(NewV, V->getType(), "", &*InsertPos));
1179  } else {
1180  U.set(ConstantExpr::getAddrSpaceCast(cast<Constant>(NewV),
1181  V->getType()));
1182  }
1183  }
1184  }
1185 
1186  if (V->use_empty()) {
1187  if (Instruction *I = dyn_cast<Instruction>(V))
1188  DeadInstructions.push_back(I);
1189  }
1190  }
1191 
1192  for (Instruction *I : DeadInstructions)
1194 
1195  return true;
1196 }
1197 
1199  if (skipFunction(F))
1200  return false;
1201 
1202  return InferAddressSpacesImpl(
1203  &getAnalysis<TargetTransformInfoWrapperPass>().getTTI(F),
1204  FlatAddrSpace)
1205  .run(F);
1206 }
1207 
1209  return new InferAddressSpaces(AddressSpace);
1210 }
1211 
1213  : FlatAddrSpace(UninitializedAddressSpace) {}
1215  : FlatAddrSpace(AddressSpace) {}
1216 
1219  bool Changed =
1220  InferAddressSpacesImpl(&AM.getResult<TargetIRAnalysis>(F), FlatAddrSpace)
1221  .run(F);
1222  if (Changed) {
1223  PreservedAnalyses PA;
1224  PA.preserveSet<CFGAnalyses>();
1225  return PA;
1226  }
1227  return PreservedAnalyses::all();
1228 }
llvm::PreservedAnalyses
A set of analyses that are preserved following a run of a transformation pass.
Definition: PassManager.h:155
llvm::RecursivelyDeleteTriviallyDeadInstructions
bool RecursivelyDeleteTriviallyDeadInstructions(Value *V, const TargetLibraryInfo *TLI=nullptr, MemorySSAUpdater *MSSAU=nullptr, std::function< void(Value *)> AboutToDeleteCallback=std::function< void(Value *)>())
If the specified value is a trivially dead instruction, delete it.
Definition: Local.cpp:511
llvm::TargetIRAnalysis
Analysis pass providing the TargetTransformInfo.
Definition: TargetTransformInfo.h:2331
llvm::GetElementPtrInst::setIsInBounds
void setIsInBounds(bool b=true)
Set or clear the inbounds flag on this GEP instruction.
Definition: Instructions.cpp:1810
MI
IRTranslator LLVM IR MI
Definition: IRTranslator.cpp:105
llvm
This file implements support for optimizing divisions by a constant.
Definition: AllocatorList.h:23
llvm::Operator
This is a utility class that provides an abstraction for the common functionality between Instruction...
Definition: Operator.h:31
M
We currently emits eax Perhaps this is what we really should generate is Is imull three or four cycles eax eax The current instruction priority is based on pattern complexity The former is more complex because it folds a load so the latter will not be emitted Perhaps we should use AddedComplexity to give LEA32r a higher priority We should always try to match LEA first since the LEA matching code does some estimate to determine whether the match is profitable if we care more about code then imull is better It s two bytes shorter than movl leal On a Pentium M
Definition: README.txt:252
Optional.h
llvm::WeakTrackingVH
Value handle that is nullable, but tries to track the Value.
Definition: ValueHandle.h:204
ValueMapper.h
llvm::CastInst::isNoopCast
static bool isNoopCast(Instruction::CastOps Opcode, Type *SrcTy, Type *DstTy, const DataLayout &DL)
A no-op cast is one that can be effected without changing any bits.
Definition: Instructions.cpp:2811
llvm::DataLayout
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:113
llvm::Intrinsic::getDeclaration
Function * getDeclaration(Module *M, ID id, ArrayRef< Type * > Tys=None)
Create or insert an LLVM Function declaration for an intrinsic, and return it.
Definition: Function.cpp:1379
llvm::BasicBlock::iterator
InstListType::iterator iterator
Instruction iterators...
Definition: BasicBlock.h:90
llvm::BasicBlock::getParent
const Function * getParent() const
Return the enclosing method, or null if none.
Definition: BasicBlock.h:107
IntrinsicInst.h
llvm::Type::isPointerTy
bool isPointerTy() const
True if this is an instance of PointerType.
Definition: Type.h:217
llvm::AnalysisManager::getResult
PassT::Result & getResult(IRUnitT &IR, ExtraArgTs... ExtraArgs)
Get the result of an analysis pass for a given IR unit.
Definition: PassManager.h:783
Scalar.h
InstIterator.h
llvm::Function
Definition: Function.h:62
Pass.h
llvm::IntrinsicInst::getIntrinsicID
Intrinsic::ID getIntrinsicID() const
Return the intrinsic ID of this intrinsic.
Definition: IntrinsicInst.h:52
llvm::CallBase::setCalledFunction
void setCalledFunction(Function *Fn)
Sets the function called, including updating the function type.
Definition: InstrTypes.h:1419
InferAddressSpaces.h
llvm::BitCastInst
This class represents a no-op cast from one type to another.
Definition: Instructions.h:5198
INITIALIZE_PASS
INITIALIZE_PASS(InferAddressSpaces, DEBUG_TYPE, "Infer address spaces", false, false) static bool isNoopPtrIntCastPair(const Operator *I2P
C1
instcombine should handle this C2 when C1
Definition: README.txt:263
llvm::SmallVector
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1168
llvm::initializeInferAddressSpacesPass
void initializeInferAddressSpacesPass(PassRegistry &)
llvm::PHINode::getOperandNumForIncomingValue
static unsigned getOperandNumForIncomingValue(unsigned i)
Definition: Instructions.h:2739
ErrorHandling.h
llvm::TargetTransformInfo
This pass provides access to the codegen interfaces that are needed for IR-level transformations.
Definition: TargetTransformInfo.h:168
llvm::Type::getPointerAddressSpace
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
Definition: DerivedTypes.h:734
llvm::IRBuilder<>
llvm::Use::get
Value * get() const
Definition: Use.h:67
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:3076
llvm::ConstantExpr::getBitCast
static Constant * getBitCast(Constant *C, Type *Ty, bool OnlyIfReduced=false)
Definition: Constants.cpp:2233
llvm::TargetTransformInfo::getAssumedAddrSpace
unsigned getAssumedAddrSpace(const Value *V) const
Definition: TargetTransformInfo.cpp:267
Local.h
llvm::ValueMap::empty
bool empty() const
Definition: ValueMap.h:140
llvm::ConstantExpr::getSelect
static Constant * getSelect(Constant *C, Constant *V1, Constant *V2, Type *OnlyIfReducedTy=nullptr)
Select constant expr.
Definition: Constants.cpp:2420
llvm::CallBase::getArgOperandUse
const Use & getArgOperandUse(unsigned i) const
Wrappers for getting the Use of a call argument.
Definition: InstrTypes.h:1339
llvm::Type
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
DenseMap.h
llvm::AtomicRMWInst::getPointerOperandIndex
static unsigned getPointerOperandIndex()
Definition: Instructions.h:867
llvm::LoadInst::getPointerOperandIndex
static unsigned getPointerOperandIndex()
Definition: Instructions.h:269
llvm::Use::getOperandNo
unsigned getOperandNo() const
Return the operand # of this use in its User.
Definition: Use.cpp:33
llvm::MemIntrinsic
This is the common base class for memset/memcpy/memmove.
Definition: IntrinsicInst.h:874
llvm::Optional< unsigned >
Operator.h
llvm::detail::DenseSetImpl::insert
std::pair< iterator, bool > insert(const ValueT &V)
Definition: DenseSet.h:206
llvm::cl::ReallyHidden
@ ReallyHidden
Definition: CommandLine.h:144
Use.h
LLVM_DEBUG
#define LLVM_DEBUG(X)
Definition: Debug.h:101
F
#define F(x, y, z)
Definition: MD5.cpp:56
operandWithNewAddressSpaceOrCreateUndef
static Value * operandWithNewAddressSpaceOrCreateUndef(const Use &OperandUse, unsigned NewAddrSpace, const ValueToValueMapTy &ValueWithNewAddrSpace, SmallVectorImpl< const Use * > *UndefUsesToFix)
Definition: InferAddressSpaces.cpp:505
llvm::Optional::hasValue
constexpr bool hasValue() const
Definition: Optional.h:288
llvm::dbgs
raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition: Debug.cpp:163
Instruction.h
CommandLine.h
P2I
auto * P2I
Definition: InferAddressSpaces.cpp:244
llvm::Intrinsic::getType
FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > Tys=None)
Return the function type for an intrinsic.
Definition: Function.cpp:1335
llvm::createInferAddressSpacesPass
FunctionPass * createInferAddressSpacesPass(unsigned AddressSpace=~0u)
Definition: InferAddressSpaces.cpp:1208
llvm::SelectInst::Create
static SelectInst * Create(Value *C, Value *S1, Value *S2, const Twine &NameStr="", Instruction *InsertBefore=nullptr, Instruction *MDFrom=nullptr)
Definition: Instructions.h:1769
llvm::AddrSpaceCastInst
This class represents a conversion between pointers from one address space to another.
Definition: Instructions.h:5238
Constants.h
llvm::Value::use_iterator
use_iterator_impl< Use > use_iterator
Definition: Value.h:354
E
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
llvm::User
Definition: User.h:44
Intrinsics.h
C
(vector float) vec_cmpeq(*A, *B) C
Definition: README_ALTIVEC.txt:86
llvm::AnalysisUsage
Represent the analysis usage information of a pass.
Definition: PassAnalysisSupport.h:47
cloneConstantExprWithNewAddressSpace
static Value * cloneConstantExprWithNewAddressSpace(ConstantExpr *CE, unsigned NewAddrSpace, const ValueToValueMapTy &ValueWithNewAddrSpace, const DataLayout *DL, const TargetTransformInfo *TTI)
Definition: InferAddressSpaces.cpp:634
DenseSet.h
llvm::MaybeAlign
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
Definition: Alignment.h:109
B
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
llvm::Instruction::CastOps
CastOps
Definition: Instruction.h:799
llvm::Instruction
Definition: Instruction.h:45
isAddressExpression
return static CastInst::isNoopCast(Instruction::CastOps(I2P->getOpcode()), I2P->getOperand(0) ->getType(), I2P->getType(), DL) &&CastInst bool isAddressExpression(const Value &V, const DataLayout &DL, const TargetTransformInfo *TTI)
Definition: InferAddressSpaces.cpp:275
llvm::PassRegistry
PassRegistry - This class manages the registration and intitialization of the pass subsystem as appli...
Definition: PassRegistry.h:38
llvm::UndefValue::get
static UndefValue * get(Type *T)
Static factory methods - Return an 'undef' object of the specified type.
Definition: Constants.cpp:1796
llvm::Use::getUser
User * getUser() const
Returns the User that contains this Use.
Definition: Use.h:73
llvm::MCID::Call
@ Call
Definition: MCInstrDesc.h:153
llvm::AddressSpace
AddressSpace
Definition: NVPTXBaseInfo.h:21
llvm::PHINode::getNumIncomingValues
unsigned getNumIncomingValues() const
Return the number of incoming edges.
Definition: Instructions.h:2725
llvm::None
const NoneType None
Definition: None.h:23
llvm::Value::use_empty
bool use_empty() const
Definition: Value.h:345
Type.h
llvm::TargetTransformInfo::hasVolatileVariant
bool hasVolatileVariant(Instruction *I, unsigned AddrSpace) const
Return true if the given instruction (assumed to be a memory access instruction) has a volatile varia...
Definition: TargetTransformInfo.cpp:422
DEBUG_TYPE
#define DEBUG_TYPE
Definition: InferAddressSpaces.cpp:134
llvm::TargetTransformInfo::getFlatAddressSpace
unsigned getFlatAddressSpace() const
Returns the address space ID for a target's 'flat' address space.
Definition: TargetTransformInfo.cpp:248
llvm::DenseSet< Value * >
llvm::Use::set
void set(Value *Val)
Definition: Value.h:864
BasicBlock.h
llvm::cl::opt< bool >
llvm::instructions
inst_range instructions(Function *F)
Definition: InstIterator.h:133
llvm::Constant
This is an important base class in LLVM.
Definition: Constant.h:41
llvm::ICmpInst
This instruction compares its operands according to the predicate given to the constructor.
Definition: Instructions.h:1203
Index
uint32_t Index
Definition: ELFObjHandler.cpp:84
llvm::TargetTransformInfo::collectFlatAddressOperands
bool collectFlatAddressOperands(SmallVectorImpl< int > &OpIndexes, Intrinsic::ID IID) const
Return any intrinsic address operand indexes which may be rewritten if they use a flat address space ...
Definition: TargetTransformInfo.cpp:252
llvm::TargetTransformInfoWrapperPass
Wrapper pass for TargetTransformInfo.
Definition: TargetTransformInfo.h:2387
llvm::GlobalValue::getParent
Module * getParent()
Get the module that this global value is contained inside of...
Definition: GlobalValue.h:578
const
aarch64 promote const
Definition: AArch64PromoteConstant.cpp:232
llvm::ConstantExpr::getAddrSpaceCast
static Constant * getAddrSpaceCast(Constant *C, Type *Ty, bool OnlyIfReduced=false)
Definition: Constants.cpp:2245
llvm::PHINode::addIncoming
void addIncoming(Value *V, BasicBlock *BB)
Add an incoming value to the end of the PHI list.
Definition: Instructions.h:2783
llvm::DenseMap< const Value *, unsigned >
I
#define I(x, y, z)
Definition: MD5.cpp:59
llvm::GetElementPtrInst
an instruction for type-safe pointer arithmetic to access elements of arrays and structs
Definition: Instructions.h:928
llvm::cl::init
initializer< Ty > init(const Ty &Val)
Definition: CommandLine.h:441
llvm::AtomicCmpXchgInst::getPointerOperandIndex
static unsigned getPointerOperandIndex()
Definition: Instructions.h:653
ArrayRef.h
llvm::TargetTransformInfo::rewriteIntrinsicWithAddressSpace
Value * rewriteIntrinsicWithAddressSpace(IntrinsicInst *II, Value *OldV, Value *NewV) const
Rewrite intrinsic call II such that OldV will be replaced with NewV, which has a different address sp...
Definition: TargetTransformInfo.cpp:271
IRBuilder.h
assert
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
SI
StandardInstrumentations SI(Debug, VerifyEach)
llvm::Value::use_begin
use_iterator use_begin()
Definition: Value.h:361
llvm::Module
A Module instance is used to store all the information related to an LLVM module.
Definition: Module.h:67
llvm::MDNode
Metadata node.
Definition: Metadata.h:906
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:954
llvm::ArrayRef
ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...
Definition: APInt.h:32
None.h
llvm::AnalysisUsage::setPreservesCFG
void setPreservesCFG()
This function should be called by the pass, iff they do not:
Definition: Pass.cpp:253
llvm_unreachable
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
Definition: ErrorHandling.h:134
llvm::Value::getType
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:256
llvm::CFGAnalyses
Represents analyses that only rely on functions' control flow.
Definition: PassManager.h:116
llvm::Value::replaceAllUsesWith
void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
Definition: Value.cpp:532
llvm::ms_demangle::IntrinsicFunctionKind::New
@ New
Compiler.h
llvm::Value::use_end
use_iterator use_end()
Definition: Value.h:369
llvm::ValueMap< const Value *, WeakTrackingVH >
ValueHandle.h
skipToNextUser
static Value::use_iterator skipToNextUser(Value::use_iterator I, Value::use_iterator End)
Definition: InferAddressSpaces.cpp:1016
llvm::CallBase::setArgOperand
void setArgOperand(unsigned i, Value *v)
Definition: InstrTypes.h:1333
llvm::MCID::Select
@ Select
Definition: MCInstrDesc.h:162
runOnFunction
static bool runOnFunction(Function &F, bool PostInlining)
Definition: EntryExitInstrumenter.cpp:69
Constant.h
llvm::ConstantExpr
A constant value that is initialized with an expression using other constant values.
Definition: Constants.h:936
getPointerOperands
static SmallVector< Value *, 2 > getPointerOperands(const Value &V, const DataLayout &DL, const TargetTransformInfo *TTI)
Definition: InferAddressSpaces.cpp:307
llvm::AMDGPU::SendMsg::Op
Op
Definition: SIDefines.h:324
llvm::PreservedAnalyses::all
static PreservedAnalyses all()
Construct a special preserved set that preserves all passes.
Definition: PassManager.h:161
llvm::PHINode::Create
static PHINode * Create(Type *Ty, unsigned NumReservedValues, const Twine &NameStr="", Instruction *InsertBefore=nullptr)
Constructors - NumReservedValues is a hint for the number of incoming edges that this phi node will h...
Definition: Instructions.h:2675
llvm::ArrayRef::begin
iterator begin() const
Definition: ArrayRef.h:153
llvm::X86::FirstMacroFusionInstKind::Cmp
@ Cmp
DL
const DataLayout & DL
Definition: InferAddressSpaces.cpp:241
llvm::TargetTransformInfo::isNoopAddrSpaceCast
bool isNoopAddrSpaceCast(unsigned FromAS, unsigned ToAS) const
Definition: TargetTransformInfo.cpp:257
Casting.h
Function.h
PassManager.h
isSimplePointerUseValidToReplace
static bool isSimplePointerUseValidToReplace(const TargetTransformInfo &TTI, Use &U, unsigned AddrSpace)
returns true if U is the pointer operand of a memory instruction with a single pointer operand that c...
Definition: InferAddressSpaces.cpp:906
llvm::IntrinsicInst
A wrapper class for inspecting calls to intrinsic functions.
Definition: IntrinsicInst.h:45
Instructions.h
llvm::PreservedAnalyses::preserveSet
void preserveSet()
Mark an analysis set as preserved.
Definition: PassManager.h:191
SmallVector.h
User.h
handleMemIntrinsicPtrUse
static bool handleMemIntrinsicPtrUse(MemIntrinsic *MI, Value *OldV, Value *NewV)
Update memory intrinsic uses that require more complex processing than simple memory instructions.
Definition: InferAddressSpaces.cpp:936
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 is ...
Definition: DerivedTypes.h:666
llvm::Type::getPointerElementType
Type * getPointerElementType() const
Definition: Type.h:369
llvm::CallBase::getArgOperand
Value * getArgOperand(unsigned i) const
Definition: InstrTypes.h:1328
llvm::Instruction::getParent
const BasicBlock * getParent() const
Definition: Instruction.h:94
llvm::PHINode::getIncomingBlock
BasicBlock * getIncomingBlock(unsigned i) const
Return incoming basic block number i.
Definition: Instructions.h:2749
llvm::max
Align max(MaybeAlign Lhs, Align Rhs)
Definition: Alignment.h:340
TargetTransformInfo.h
llvm::PHINode
Definition: Instructions.h:2633
llvm::SmallVectorImpl
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
Definition: APFloat.h:43
llvm::AnalysisManager
A container for analyses that lazily runs them and caches their results.
Definition: InstructionSimplify.h:44
AssumeDefaultIsFlatAddressSpace
static cl::opt< bool > AssumeDefaultIsFlatAddressSpace("assume-default-is-flat-addrspace", cl::init(false), cl::ReallyHidden, cl::desc("The default address space is assumed as the flat address space. " "This is mainly for test purpose."))
llvm::FunctionPass
FunctionPass class - This class is used to implement most global optimizations.
Definition: Pass.h:298
UninitializedAddressSpace
static const unsigned UninitializedAddressSpace
Definition: InferAddressSpaces.cpp:143
GEP
Hexagon Common GEP
Definition: HexagonCommonGEP.cpp:172
llvm::ValueMap::lookup
ValueT lookup(const KeyT &Val) const
lookup - Return the entry for the specified key, or a default constructed value if no such entry exis...
Definition: ValueMap.h:165
llvm::AnalysisUsage::addRequired
AnalysisUsage & addRequired()
Definition: PassAnalysisSupport.h:75
LLVMContext.h
llvm::User::getOperand
Value * getOperand(unsigned i) const
Definition: User.h:169
llvm::cl::desc
Definition: CommandLine.h:412
raw_ostream.h
llvm::StoreInst::getPointerOperandIndex
static unsigned getPointerOperandIndex()
Definition: Instructions.h:403
llvm::SetVector< Value * >
Value.h
llvm::Value
LLVM Value Representation.
Definition: Value.h:75
llvm::InferAddressSpacesPass::run
PreservedAnalyses run(Function &F, FunctionAnalysisManager &AM)
Definition: InferAddressSpaces.cpp:1217
llvm::InferAddressSpacesPass::InferAddressSpacesPass
InferAddressSpacesPass()
Definition: InferAddressSpaces.cpp:1212
Debug.h
llvm::Value::users
iterator_range< user_iterator > users()
Definition: Value.h:422
llvm::ArrayRef::end
iterator end() const
Definition: ArrayRef.h:154
llvm::Optional::getValue
constexpr const T & getValue() const LLVM_LVALUE_FUNCTION
Definition: Optional.h:282
SetVector.h
llvm::Use
A Use represents the edge between a Value definition and its users.
Definition: Use.h:44
llvm::Intrinsic::ID
unsigned ID
Definition: TargetTransformInfo.h:37