LLVM  14.0.0git
AMDGPULDSUtils.cpp
Go to the documentation of this file.
1 //===- AMDGPULDSUtils.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 // AMDGPU LDS related helper utility functions.
10 //
11 //===----------------------------------------------------------------------===//
12 
13 #include "AMDGPULDSUtils.h"
14 #include "AMDGPU.h"
15 #include "Utils/AMDGPUBaseInfo.h"
17 #include "llvm/ADT/SetVector.h"
18 #include "llvm/IR/Constants.h"
20 
21 using namespace llvm;
22 
23 namespace llvm {
24 
25 namespace AMDGPU {
26 
27 bool isKernelCC(const Function *Func) {
28  return AMDGPU::isModuleEntryFunctionCC(Func->getCallingConv());
29 }
30 
32  return DL.getValueOrABITypeAlignment(GV->getPointerAlignment(DL),
33  GV->getValueType());
34 }
35 
36 static void collectFunctionUses(User *U, const Function *F,
37  SetVector<Instruction *> &InstUsers) {
38  SmallVector<User *> Stack{U};
39 
40  while (!Stack.empty()) {
41  U = Stack.pop_back_val();
42 
43  if (auto *I = dyn_cast<Instruction>(U)) {
44  if (I->getFunction() == F)
45  InstUsers.insert(I);
46  continue;
47  }
48 
49  if (!isa<ConstantExpr>(U))
50  continue;
51 
52  append_range(Stack, U->users());
53  }
54 }
55 
57  SetVector<Instruction *> InstUsers;
58 
59  collectFunctionUses(C, F, InstUsers);
60  for (Instruction *I : InstUsers) {
62  }
63 }
64 
65 static bool shouldLowerLDSToStruct(const GlobalVariable &GV,
66  const Function *F) {
67  // We are not interested in kernel LDS lowering for module LDS itself.
68  if (F && GV.getName() == "llvm.amdgcn.module.lds")
69  return false;
70 
71  bool Ret = false;
74 
75  assert(!F || isKernelCC(F));
76 
77  while (!Stack.empty()) {
78  const User *V = Stack.pop_back_val();
79  Visited.insert(V);
80 
81  if (isa<GlobalValue>(V)) {
82  // This use of the LDS variable is the initializer of a global variable.
83  // This is ill formed. The address of an LDS variable is kernel dependent
84  // and unknown until runtime. It can't be written to a global variable.
85  continue;
86  }
87 
88  if (auto *I = dyn_cast<Instruction>(V)) {
89  const Function *UF = I->getFunction();
90  if (UF == F) {
91  // Used from this kernel, we want to put it into the structure.
92  Ret = true;
93  } else if (!F) {
94  // For module LDS lowering, lowering is required if the user instruction
95  // is from non-kernel function.
96  Ret |= !isKernelCC(UF);
97  }
98  continue;
99  }
100 
101  // User V should be a constant, recursively visit users of V.
102  assert(isa<Constant>(V) && "Expected a constant.");
103  append_range(Stack, V->users());
104  }
105 
106  return Ret;
107 }
108 
109 std::vector<GlobalVariable *> findVariablesToLower(Module &M,
110  const Function *F) {
111  std::vector<llvm::GlobalVariable *> LocalVars;
112  for (auto &GV : M.globals()) {
113  if (GV.getType()->getPointerAddressSpace() != AMDGPUAS::LOCAL_ADDRESS) {
114  continue;
115  }
116  if (!GV.hasInitializer()) {
117  // addrspace(3) without initializer implies cuda/hip extern __shared__
118  // the semantics for such a variable appears to be that all extern
119  // __shared__ variables alias one another, in which case this transform
120  // is not required
121  continue;
122  }
123  if (!isa<UndefValue>(GV.getInitializer())) {
124  // Initializers are unimplemented for LDS address space.
125  // Leave such variables in place for consistent error reporting.
126  continue;
127  }
128  if (GV.isConstant()) {
129  // A constant undef variable can't be written to, and any load is
130  // undef, so it should be eliminated by the optimizer. It could be
131  // dropped by the back end if not. This pass skips over it.
132  continue;
133  }
134  if (!shouldLowerLDSToStruct(GV, F)) {
135  continue;
136  }
137  LocalVars.push_back(&GV);
138  }
139  return LocalVars;
140 }
141 
142 } // end namespace AMDGPU
143 
144 } // end namespace llvm
llvm::AMDGPU::findVariablesToLower
std::vector< GlobalVariable * > findVariablesToLower(Module &M, const Function *F)
Definition: AMDGPULDSUtils.cpp:109
llvm
This is an optimization pass for GlobalISel generic memory operations.
Definition: AllocatorList.h:23
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
llvm::Value::getPointerAlignment
Align getPointerAlignment(const DataLayout &DL) const
Returns an alignment of the pointer value.
Definition: Value.cpp:918
llvm::DataLayout
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:113
llvm::Function
Definition: Function.h:62
llvm::SmallVector
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1175
llvm::GlobalVariable
Definition: GlobalVariable.h:39
llvm::SmallPtrSet
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
Definition: SmallPtrSet.h:449
llvm::AMDGPU::shouldLowerLDSToStruct
static bool shouldLowerLDSToStruct(const GlobalVariable &GV, const Function *F)
Definition: AMDGPULDSUtils.cpp:65
llvm::MipsISD::Ret
@ Ret
Definition: MipsISelLowering.h:116
DepthFirstIterator.h
F
#define F(x, y, z)
Definition: MD5.cpp:55
AMDGPULDSUtils.h
Constants.h
llvm::User
Definition: User.h:44
C
(vector float) vec_cmpeq(*A, *B) C
Definition: README_ALTIVEC.txt:86
llvm::convertConstantExprsToInstructions
void convertConstantExprsToInstructions(Instruction *I, ConstantExpr *CE, SmallPtrSetImpl< Instruction * > *Insts=nullptr)
The given instruction I contains given constant expression CE as one of its operands,...
Definition: ReplaceConstant.cpp:22
AMDGPU
Definition: AMDGPUReplaceLDSUseWithPointer.cpp:114
llvm::Instruction
Definition: Instruction.h:45
llvm::Align
This struct is a compact representation of a valid (non-zero power of two) alignment.
Definition: Alignment.h:39
llvm::AMDGPU::collectFunctionUses
static void collectFunctionUses(User *U, const Function *F, SetVector< Instruction * > &InstUsers)
Definition: AMDGPULDSUtils.cpp:36
llvm::AMDGPU::replaceConstantUsesInFunction
void replaceConstantUsesInFunction(ConstantExpr *C, const Function *F)
Replace all uses of constant C with instructions in F.
Definition: AMDGPULDSUtils.cpp:56
I
#define I(x, y, z)
Definition: MD5.cpp:58
assert
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
llvm::Module
A Module instance is used to store all the information related to an LLVM module.
Definition: Module.h:65
llvm::SetVector::insert
bool insert(const value_type &X)
Insert a new element into the SetVector.
Definition: SetVector.h:141
AMDGPU.h
llvm::AMDGPU::isModuleEntryFunctionCC
bool isModuleEntryFunctionCC(CallingConv::ID CC)
Definition: AMDGPUBaseInfo.cpp:1411
llvm::append_range
void append_range(Container &C, Range &&R)
Wrapper function to append a range to a container.
Definition: STLExtras.h:1789
DL
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
Definition: AArch64SLSHardening.cpp:76
llvm::Value::getName
StringRef getName() const
Return a constant reference to the value's name.
Definition: Value.cpp:309
llvm::AMDGPU::getAlign
Align getAlign(DataLayout const &DL, const GlobalVariable *GV)
Definition: AMDGPULDSUtils.cpp:31
llvm::ConstantExpr
A constant value that is initialized with an expression using other constant values.
Definition: Constants.h:971
ReplaceConstant.h
llvm::GlobalValue::getValueType
Type * getValueType() const
Definition: GlobalValue.h:273
llvm::SetVector
A vector that has set insertion semantics.
Definition: SetVector.h:40
llvm::AMDGPU::isKernelCC
bool isKernelCC(const Function *Func)
Definition: AMDGPULDSUtils.cpp:27
llvm::AMDGPUAS::LOCAL_ADDRESS
@ LOCAL_ADDRESS
Address space for local memory.
Definition: AMDGPU.h:363
llvm::Value::users
iterator_range< user_iterator > users()
Definition: Value.h:421
SetVector.h
AMDGPUBaseInfo.h
llvm::SmallPtrSetImpl::insert
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
Definition: SmallPtrSet.h:364