Bug Summary

File:llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
Warning:line 236, column 22
Called C++ object pointer is null

Annotated Source Code

Press '?' to see keyboard shortcuts

clang -cc1 -cc1 -triple x86_64-pc-linux-gnu -analyze -disable-free -disable-llvm-verifier -discard-value-names -main-file-name NVPTXLowerArgs.cpp -analyzer-store=region -analyzer-opt-analyze-nested-blocks -analyzer-checker=core -analyzer-checker=apiModeling -analyzer-checker=unix -analyzer-checker=deadcode -analyzer-checker=cplusplus -analyzer-checker=security.insecureAPI.UncheckedReturn -analyzer-checker=security.insecureAPI.getpw -analyzer-checker=security.insecureAPI.gets -analyzer-checker=security.insecureAPI.mktemp -analyzer-checker=security.insecureAPI.mkstemp -analyzer-checker=security.insecureAPI.vfork -analyzer-checker=nullability.NullPassedToNonnull -analyzer-checker=nullability.NullReturnedFromNonnull -analyzer-output plist -w -setup-static-analyzer -analyzer-config-compatibility-mode=true -mrelocation-model pic -pic-level 2 -mframe-pointer=none -fmath-errno -fno-rounding-math -mconstructor-aliases -munwind-tables -target-cpu x86-64 -tune-cpu generic -debugger-tuning=gdb -ffunction-sections -fdata-sections -fcoverage-compilation-dir=/build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/lib/Target/NVPTX -resource-dir /usr/lib/llvm-14/lib/clang/14.0.0 -D _GNU_SOURCE -D __STDC_CONSTANT_MACROS -D __STDC_FORMAT_MACROS -D __STDC_LIMIT_MACROS -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/lib/Target/NVPTX -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/llvm/lib/Target/NVPTX -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/include -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/llvm/include -D NDEBUG -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/x86_64-linux-gnu/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/backward -internal-isystem /usr/lib/llvm-14/lib/clang/14.0.0/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../x86_64-linux-gnu/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -O2 -Wno-unused-parameter -Wwrite-strings -Wno-missing-field-initializers -Wno-long-long -Wno-maybe-uninitialized -Wno-class-memaccess -Wno-redundant-move -Wno-pessimizing-move -Wno-noexcept-type -Wno-comment -std=c++14 -fdeprecated-macro -fdebug-compilation-dir=/build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/lib/Target/NVPTX -fdebug-prefix-map=/build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e=. -ferror-limit 19 -fvisibility hidden -fvisibility-inlines-hidden -stack-protector 2 -fgnuc-version=4.2.1 -vectorize-loops -vectorize-slp -analyzer-output=html -analyzer-config stable-report-filename=true -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/scan-build-2021-09-04-040900-46481-1 -x c++ /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
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
91#include "NVPTX.h"
92#include "NVPTXTargetMachine.h"
93#include "NVPTXUtilities.h"
94#include "MCTargetDesc/NVPTXBaseInfo.h"
95#include "llvm/Analysis/ValueTracking.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#define DEBUG_TYPE"nvptx-lower-args" "nvptx-lower-args"
103
104using namespace llvm;
105
106namespace llvm {
107void initializeNVPTXLowerArgsPass(PassRegistry &);
108}
109
110namespace {
111class NVPTXLowerArgs : public FunctionPass {
112 bool runOnFunction(Function &F) override;
113
114 bool runOnKernelFunction(Function &F);
115 bool runOnDeviceFunction(Function &F);
116
117 // handle byval parameters
118 void handleByValParam(Argument *Arg);
119 // Knowing Ptr must point to the global address space, this function
120 // addrspacecasts Ptr to global and then back to generic. This allows
121 // NVPTXInferAddressSpaces to fold the global-to-generic cast into
122 // loads/stores that appear later.
123 void markPointerAsGlobal(Value *Ptr);
124
125public:
126 static char ID; // Pass identification, replacement for typeid
127 NVPTXLowerArgs(const NVPTXTargetMachine *TM = nullptr)
128 : FunctionPass(ID), TM(TM) {}
129 StringRef getPassName() const override {
130 return "Lower pointer arguments of CUDA kernels";
131 }
132
133private:
134 const NVPTXTargetMachine *TM;
135};
136} // namespace
137
138char NVPTXLowerArgs::ID = 1;
139
140INITIALIZE_PASS(NVPTXLowerArgs, "nvptx-lower-args",static void *initializeNVPTXLowerArgsPassOnce(PassRegistry &
Registry) { PassInfo *PI = new PassInfo( "Lower arguments (NVPTX)"
, "nvptx-lower-args", &NVPTXLowerArgs::ID, PassInfo::NormalCtor_t
(callDefaultCtor<NVPTXLowerArgs>), false, false); Registry
.registerPass(*PI, true); return PI; } static llvm::once_flag
InitializeNVPTXLowerArgsPassFlag; void llvm::initializeNVPTXLowerArgsPass
(PassRegistry &Registry) { llvm::call_once(InitializeNVPTXLowerArgsPassFlag
, initializeNVPTXLowerArgsPassOnce, std::ref(Registry)); }
141 "Lower arguments (NVPTX)", false, false)static void *initializeNVPTXLowerArgsPassOnce(PassRegistry &
Registry) { PassInfo *PI = new PassInfo( "Lower arguments (NVPTX)"
, "nvptx-lower-args", &NVPTXLowerArgs::ID, PassInfo::NormalCtor_t
(callDefaultCtor<NVPTXLowerArgs>), false, false); Registry
.registerPass(*PI, true); return PI; } static llvm::once_flag
InitializeNVPTXLowerArgsPassFlag; void llvm::initializeNVPTXLowerArgsPass
(PassRegistry &Registry) { llvm::call_once(InitializeNVPTXLowerArgsPassFlag
, initializeNVPTXLowerArgsPassOnce, std::ref(Registry)); }
142
143// =============================================================================
144// If the function had a byval struct ptr arg, say foo(%struct.x* byval %d),
145// and we can't guarantee that the only accesses are loads,
146// then add the following instructions to the first basic block:
147//
148// %temp = alloca %struct.x, align 8
149// %tempd = addrspacecast %struct.x* %d to %struct.x addrspace(101)*
150// %tv = load %struct.x addrspace(101)* %tempd
151// store %struct.x %tv, %struct.x* %temp, align 8
152//
153// The above code allocates some space in the stack and copies the incoming
154// struct from param space to local space.
155// Then replace all occurrences of %d by %temp.
156//
157// In case we know that all users are GEPs or Loads, replace them with the same
158// ones in parameter AS, so we can access them using ld.param.
159// =============================================================================
160
161// Replaces the \p OldUser instruction with the same in parameter AS.
162// Only Load and GEP are supported.
163static void convertToParamAS(Value *OldUser, Value *Param) {
164 Instruction *I = dyn_cast<Instruction>(OldUser);
165 assert(I && "OldUser must be an instruction")(static_cast<void> (0));
166 struct IP {
167 Instruction *OldInstruction;
168 Value *NewParam;
169 };
170 SmallVector<IP> ItemsToConvert = {{I, Param}};
171 SmallVector<Instruction *> InstructionsToDelete;
172
173 auto CloneInstInParamAS = [](const IP &I) -> Value * {
174 if (auto *LI = dyn_cast<LoadInst>(I.OldInstruction)) {
175 LI->setOperand(0, I.NewParam);
176 return LI;
177 }
178 if (auto *GEP = dyn_cast<GetElementPtrInst>(I.OldInstruction)) {
179 SmallVector<Value *, 4> Indices(GEP->indices());
180 auto *NewGEP = GetElementPtrInst::Create(GEP->getSourceElementType(),
181 I.NewParam, Indices,
182 GEP->getName(), GEP);
183 NewGEP->setIsInBounds(GEP->isInBounds());
184 return NewGEP;
185 }
186 if (auto *BC = dyn_cast<BitCastInst>(I.OldInstruction)) {
187 auto *NewBCType = PointerType::getWithSamePointeeType(
188 cast<PointerType>(BC->getType()), ADDRESS_SPACE_PARAM);
189 return BitCastInst::Create(BC->getOpcode(), I.NewParam, NewBCType,
190 BC->getName(), BC);
191 }
192 if (auto *ASC = dyn_cast<AddrSpaceCastInst>(I.OldInstruction)) {
193 assert(ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM)(static_cast<void> (0));
194 (void)ASC;
195 // Just pass through the argument, the old ASC is no longer needed.
196 return I.NewParam;
197 }
198 llvm_unreachable("Unsupported instruction")__builtin_unreachable();
199 };
200
201 while (!ItemsToConvert.empty()) {
202 IP I = ItemsToConvert.pop_back_val();
203 Value *NewInst = CloneInstInParamAS(I);
204
205 if (NewInst && NewInst != I.OldInstruction) {
206 // We've created a new instruction. Queue users of the old instruction to
207 // be converted and the instruction itself to be deleted. We can't delete
208 // the old instruction yet, because it's still in use by a load somewhere.
209 llvm::for_each(
210 I.OldInstruction->users(), [NewInst, &ItemsToConvert](Value *V) {
211 ItemsToConvert.push_back({cast<Instruction>(V), NewInst});
212 });
213
214 InstructionsToDelete.push_back(I.OldInstruction);
215 }
216 }
217
218 // Now we know that all argument loads are using addresses in parameter space
219 // and we can finally remove the old instructions in generic AS. Instructions
220 // scheduled for removal should be processed in reverse order so the ones
221 // closest to the load are deleted first. Otherwise they may still be in use.
222 // E.g if we have Value = Load(BitCast(GEP(arg))), InstructionsToDelete will
223 // have {GEP,BitCast}. GEP can't be deleted first, because it's still used by
224 // the BitCast.
225 llvm::for_each(reverse(InstructionsToDelete),
226 [](Instruction *I) { I->eraseFromParent(); });
227}
228
229void NVPTXLowerArgs::handleByValParam(Argument *Arg) {
230 Function *Func = Arg->getParent();
231 Instruction *FirstInst = &(Func->getEntryBlock().front());
232 PointerType *PType = dyn_cast<PointerType>(Arg->getType());
9
←
Assuming the object is not a 'PointerType'
→
10
←
'PType' initialized to a null pointer value
→
233
234 assert(PType && "Expecting pointer type in handleByValParam")(static_cast<void> (0));
235
236 Type *StructType = PType->getElementType();
11
←
Called C++ object pointer is null
237
238 auto IsALoadChain = [&](Value *Start) {
239 SmallVector<Value *, 16> ValuesToCheck = {Start};
240 auto IsALoadChainInstr = [](Value *V) -> bool {
241 if (isa<GetElementPtrInst>(V) || isa<BitCastInst>(V) || isa<LoadInst>(V))
242 return true;
243 // ASC to param space are OK, too -- we'll just strip them.
244 if (auto *ASC = dyn_cast<AddrSpaceCastInst>(V)) {
245 if (ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM)
246 return true;
247 }
248 return false;
249 };
250
251 while (!ValuesToCheck.empty()) {
252 Value *V = ValuesToCheck.pop_back_val();
253 if (!IsALoadChainInstr(V)) {
254 LLVM_DEBUG(dbgs() << "Need a copy of " << *Arg << " because of " << *Vdo { } while (false)
255 << "\n")do { } while (false);
256 (void)Arg;
257 return false;
258 }
259 if (!isa<LoadInst>(V))
260 llvm::append_range(ValuesToCheck, V->users());
261 }
262 return true;
263 };
264
265 if (llvm::all_of(Arg->users(), IsALoadChain)) {
266 // Convert all loads and intermediate operations to use parameter AS and
267 // skip creation of a local copy of the argument.
268 SmallVector<User *, 16> UsersToUpdate(Arg->users());
269 Value *ArgInParamAS = new AddrSpaceCastInst(
270 Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(),
271 FirstInst);
272 llvm::for_each(UsersToUpdate, [ArgInParamAS](Value *V) {
273 convertToParamAS(V, ArgInParamAS);
274 });
275 LLVM_DEBUG(dbgs() << "No need to copy " << *Arg << "\n")do { } while (false);
276 return;
277 }
278
279 // Otherwise we have to create a temporary copy.
280 const DataLayout &DL = Func->getParent()->getDataLayout();
281 unsigned AS = DL.getAllocaAddrSpace();
282 AllocaInst *AllocA = new AllocaInst(StructType, AS, Arg->getName(), FirstInst);
283 // Set the alignment to alignment of the byval parameter. This is because,
284 // later load/stores assume that alignment, and we are going to replace
285 // the use of the byval parameter with this alloca instruction.
286 AllocA->setAlignment(Func->getParamAlign(Arg->getArgNo())
287 .getValueOr(DL.getPrefTypeAlign(StructType)));
288 Arg->replaceAllUsesWith(AllocA);
289
290 Value *ArgInParam = new AddrSpaceCastInst(
291 Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(),
292 FirstInst);
293 // Be sure to propagate alignment to this load; LLVM doesn't know that NVPTX
294 // addrspacecast preserves alignment. Since params are constant, this load is
295 // definitely not volatile.
296 LoadInst *LI =
297 new LoadInst(StructType, ArgInParam, Arg->getName(),
298 /*isVolatile=*/false, AllocA->getAlign(), FirstInst);
299 new StoreInst(LI, AllocA, FirstInst);
300}
301
302void NVPTXLowerArgs::markPointerAsGlobal(Value *Ptr) {
303 if (Ptr->getType()->getPointerAddressSpace() == ADDRESS_SPACE_GLOBAL)
304 return;
305
306 // Deciding where to emit the addrspacecast pair.
307 BasicBlock::iterator InsertPt;
308 if (Argument *Arg = dyn_cast<Argument>(Ptr)) {
309 // Insert at the functon entry if Ptr is an argument.
310 InsertPt = Arg->getParent()->getEntryBlock().begin();
311 } else {
312 // Insert right after Ptr if Ptr is an instruction.
313 InsertPt = ++cast<Instruction>(Ptr)->getIterator();
314 assert(InsertPt != InsertPt->getParent()->end() &&(static_cast<void> (0))
315 "We don't call this function with Ptr being a terminator.")(static_cast<void> (0));
316 }
317
318 Instruction *PtrInGlobal = new AddrSpaceCastInst(
319 Ptr,
320 PointerType::getWithSamePointeeType(cast<PointerType>(Ptr->getType()),
321 ADDRESS_SPACE_GLOBAL),
322 Ptr->getName(), &*InsertPt);
323 Value *PtrInGeneric = new AddrSpaceCastInst(PtrInGlobal, Ptr->getType(),
324 Ptr->getName(), &*InsertPt);
325 // Replace with PtrInGeneric all uses of Ptr except PtrInGlobal.
326 Ptr->replaceAllUsesWith(PtrInGeneric);
327 PtrInGlobal->setOperand(0, Ptr);
328}
329
330// =============================================================================
331// Main function for this pass.
332// =============================================================================
333bool NVPTXLowerArgs::runOnKernelFunction(Function &F) {
334 if (TM && TM->getDrvInterface() == NVPTX::CUDA) {
335 // Mark pointers in byval structs as global.
336 for (auto &B : F) {
337 for (auto &I : B) {
338 if (LoadInst *LI = dyn_cast<LoadInst>(&I)) {
339 if (LI->getType()->isPointerTy()) {
340 Value *UO = getUnderlyingObject(LI->getPointerOperand());
341 if (Argument *Arg = dyn_cast<Argument>(UO)) {
342 if (Arg->hasByValAttr()) {
343 // LI is a load from a pointer within a byval kernel parameter.
344 markPointerAsGlobal(LI);
345 }
346 }
347 }
348 }
349 }
350 }
351 }
352
353 LLVM_DEBUG(dbgs() << "Lowering kernel args of " << F.getName() << "\n")do { } while (false);
354 for (Argument &Arg : F.args()) {
355 if (Arg.getType()->isPointerTy()) {
356 if (Arg.hasByValAttr())
357 handleByValParam(&Arg);
358 else if (TM && TM->getDrvInterface() == NVPTX::CUDA)
359 markPointerAsGlobal(&Arg);
360 }
361 }
362 return true;
363}
364
365// Device functions only need to copy byval args into local memory.
366bool NVPTXLowerArgs::runOnDeviceFunction(Function &F) {
367 LLVM_DEBUG(dbgs() << "Lowering function args of " << F.getName() << "\n")do { } while (false);
4
←
Loop condition is false. Exiting loop
→
368 for (Argument &Arg : F.args())
5
←
Assuming '__begin1' is not equal to '__end1'
→
369 if (Arg.getType()->isPointerTy() && Arg.hasByValAttr())
6
←
Assuming the condition is true
→
7
←
Taking true branch
→
370 handleByValParam(&Arg);
8
←
Calling 'NVPTXLowerArgs::handleByValParam'
→
371 return true;
372}
373
374bool NVPTXLowerArgs::runOnFunction(Function &F) {
375 return isKernelFunction(F) ? runOnKernelFunction(F) : runOnDeviceFunction(F);
1
Assuming the condition is false
→
2
←
'?' condition is false
→
3
←
Calling 'NVPTXLowerArgs::runOnDeviceFunction'
→
376}
377
378FunctionPass *
379llvm::createNVPTXLowerArgsPass(const NVPTXTargetMachine *TM) {
380 return new NVPTXLowerArgs(TM);
381}