Bug Summary

File:clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
Warning:line 4437, column 8
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 CGOpenMPRuntimeGPU.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 -fhalf-no-semantic-interposition -mframe-pointer=none -relaxed-aliasing -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-13~++20210413100635+64c24f493e5f/build-llvm/tools/clang/lib/CodeGen -resource-dir /usr/lib/llvm-13/lib/clang/13.0.0 -D CLANG_ROUND_TRIP_CC1_ARGS=ON -D _DEBUG -D _GNU_SOURCE -D __STDC_CONSTANT_MACROS -D __STDC_FORMAT_MACROS -D __STDC_LIMIT_MACROS -I /build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/build-llvm/tools/clang/lib/CodeGen -I /build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen -I /build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/include -I /build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/build-llvm/tools/clang/include -I /build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/build-llvm/include -I /build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/llvm/include -U NDEBUG -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../include/c++/6.3.0 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../include/x86_64-linux-gnu/c++/6.3.0 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../include/c++/6.3.0/backward -internal-isystem /usr/lib/llvm-13/lib/clang/13.0.0/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../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-comment -std=c++14 -fdeprecated-macro -fdebug-compilation-dir=/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/build-llvm/tools/clang/lib/CodeGen -fdebug-prefix-map=/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f=. -ferror-limit 19 -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-04-14-063029-18377-1 -x c++ /build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
1//===---- CGOpenMPRuntimeGPU.cpp - Interface to OpenMP GPU Runtimes ----===//
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// This provides a generalized class for OpenMP runtime code generation
10// specialized by GPU targets NVPTX and AMDGCN.
11//
12//===----------------------------------------------------------------------===//
13
14#include "CGOpenMPRuntimeGPU.h"
15#include "CGOpenMPRuntimeNVPTX.h"
16#include "CodeGenFunction.h"
17#include "clang/AST/Attr.h"
18#include "clang/AST/DeclOpenMP.h"
19#include "clang/AST/StmtOpenMP.h"
20#include "clang/AST/StmtVisitor.h"
21#include "clang/Basic/Cuda.h"
22#include "llvm/ADT/SmallPtrSet.h"
23#include "llvm/Frontend/OpenMP/OMPGridValues.h"
24#include "llvm/IR/IntrinsicsNVPTX.h"
25
26using namespace clang;
27using namespace CodeGen;
28using namespace llvm::omp;
29
30namespace {
31/// Pre(post)-action for different OpenMP constructs specialized for NVPTX.
32class NVPTXActionTy final : public PrePostActionTy {
33 llvm::FunctionCallee EnterCallee = nullptr;
34 ArrayRef<llvm::Value *> EnterArgs;
35 llvm::FunctionCallee ExitCallee = nullptr;
36 ArrayRef<llvm::Value *> ExitArgs;
37 bool Conditional = false;
38 llvm::BasicBlock *ContBlock = nullptr;
39
40public:
41 NVPTXActionTy(llvm::FunctionCallee EnterCallee,
42 ArrayRef<llvm::Value *> EnterArgs,
43 llvm::FunctionCallee ExitCallee,
44 ArrayRef<llvm::Value *> ExitArgs, bool Conditional = false)
45 : EnterCallee(EnterCallee), EnterArgs(EnterArgs), ExitCallee(ExitCallee),
46 ExitArgs(ExitArgs), Conditional(Conditional) {}
47 void Enter(CodeGenFunction &CGF) override {
48 llvm::Value *EnterRes = CGF.EmitRuntimeCall(EnterCallee, EnterArgs);
49 if (Conditional) {
50 llvm::Value *CallBool = CGF.Builder.CreateIsNotNull(EnterRes);
51 auto *ThenBlock = CGF.createBasicBlock("omp_if.then");
52 ContBlock = CGF.createBasicBlock("omp_if.end");
53 // Generate the branch (If-stmt)
54 CGF.Builder.CreateCondBr(CallBool, ThenBlock, ContBlock);
55 CGF.EmitBlock(ThenBlock);
56 }
57 }
58 void Done(CodeGenFunction &CGF) {
59 // Emit the rest of blocks/branches
60 CGF.EmitBranch(ContBlock);
61 CGF.EmitBlock(ContBlock, true);
62 }
63 void Exit(CodeGenFunction &CGF) override {
64 CGF.EmitRuntimeCall(ExitCallee, ExitArgs);
65 }
66};
67
68/// A class to track the execution mode when codegening directives within
69/// a target region. The appropriate mode (SPMD|NON-SPMD) is set on entry
70/// to the target region and used by containing directives such as 'parallel'
71/// to emit optimized code.
72class ExecutionRuntimeModesRAII {
73private:
74 CGOpenMPRuntimeGPU::ExecutionMode SavedExecMode =
75 CGOpenMPRuntimeGPU::EM_Unknown;
76 CGOpenMPRuntimeGPU::ExecutionMode &ExecMode;
77 bool SavedRuntimeMode = false;
78 bool *RuntimeMode = nullptr;
79
80public:
81 /// Constructor for Non-SPMD mode.
82 ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode)
83 : ExecMode(ExecMode) {
84 SavedExecMode = ExecMode;
85 ExecMode = CGOpenMPRuntimeGPU::EM_NonSPMD;
86 }
87 /// Constructor for SPMD mode.
88 ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode,
89 bool &RuntimeMode, bool FullRuntimeMode)
90 : ExecMode(ExecMode), RuntimeMode(&RuntimeMode) {
91 SavedExecMode = ExecMode;
92 SavedRuntimeMode = RuntimeMode;
93 ExecMode = CGOpenMPRuntimeGPU::EM_SPMD;
94 RuntimeMode = FullRuntimeMode;
95 }
96 ~ExecutionRuntimeModesRAII() {
97 ExecMode = SavedExecMode;
98 if (RuntimeMode)
99 *RuntimeMode = SavedRuntimeMode;
100 }
101};
102
103/// GPU Configuration: This information can be derived from cuda registers,
104/// however, providing compile time constants helps generate more efficient
105/// code. For all practical purposes this is fine because the configuration
106/// is the same for all known NVPTX architectures.
107enum MachineConfiguration : unsigned {
108 /// See "llvm/Frontend/OpenMP/OMPGridValues.h" for various related target
109 /// specific Grid Values like GV_Warp_Size, GV_Warp_Size_Log2,
110 /// and GV_Warp_Size_Log2_Mask.
111
112 /// Global memory alignment for performance.
113 GlobalMemoryAlignment = 128,
114
115 /// Maximal size of the shared memory buffer.
116 SharedMemorySize = 128,
117};
118
119static const ValueDecl *getPrivateItem(const Expr *RefExpr) {
120 RefExpr = RefExpr->IgnoreParens();
121 if (const auto *ASE = dyn_cast<ArraySubscriptExpr>(RefExpr)) {
122 const Expr *Base = ASE->getBase()->IgnoreParenImpCasts();
123 while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))
124 Base = TempASE->getBase()->IgnoreParenImpCasts();
125 RefExpr = Base;
126 } else if (auto *OASE = dyn_cast<OMPArraySectionExpr>(RefExpr)) {
127 const Expr *Base = OASE->getBase()->IgnoreParenImpCasts();
128 while (const auto *TempOASE = dyn_cast<OMPArraySectionExpr>(Base))
129 Base = TempOASE->getBase()->IgnoreParenImpCasts();
130 while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))
131 Base = TempASE->getBase()->IgnoreParenImpCasts();
132 RefExpr = Base;
133 }
134 RefExpr = RefExpr->IgnoreParenImpCasts();
135 if (const auto *DE = dyn_cast<DeclRefExpr>(RefExpr))
136 return cast<ValueDecl>(DE->getDecl()->getCanonicalDecl());
137 const auto *ME = cast<MemberExpr>(RefExpr);
138 return cast<ValueDecl>(ME->getMemberDecl()->getCanonicalDecl());
139}
140
141
142static RecordDecl *buildRecordForGlobalizedVars(
143 ASTContext &C, ArrayRef<const ValueDecl *> EscapedDecls,
144 ArrayRef<const ValueDecl *> EscapedDeclsForTeams,
145 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
146 &MappedDeclsFields, int BufSize) {
147 using VarsDataTy = std::pair<CharUnits /*Align*/, const ValueDecl *>;
148 if (EscapedDecls.empty() && EscapedDeclsForTeams.empty())
149 return nullptr;
150 SmallVector<VarsDataTy, 4> GlobalizedVars;
151 for (const ValueDecl *D : EscapedDecls)
152 GlobalizedVars.emplace_back(
153 CharUnits::fromQuantity(std::max(
154 C.getDeclAlign(D).getQuantity(),
155 static_cast<CharUnits::QuantityType>(GlobalMemoryAlignment))),
156 D);
157 for (const ValueDecl *D : EscapedDeclsForTeams)
158 GlobalizedVars.emplace_back(C.getDeclAlign(D), D);
159 llvm::stable_sort(GlobalizedVars, [](VarsDataTy L, VarsDataTy R) {
160 return L.first > R.first;
161 });
162
163 // Build struct _globalized_locals_ty {
164 // /* globalized vars */[WarSize] align (max(decl_align,
165 // GlobalMemoryAlignment))
166 // /* globalized vars */ for EscapedDeclsForTeams
167 // };
168 RecordDecl *GlobalizedRD = C.buildImplicitRecord("_globalized_locals_ty");
169 GlobalizedRD->startDefinition();
170 llvm::SmallPtrSet<const ValueDecl *, 16> SingleEscaped(
171 EscapedDeclsForTeams.begin(), EscapedDeclsForTeams.end());
172 for (const auto &Pair : GlobalizedVars) {
173 const ValueDecl *VD = Pair.second;
174 QualType Type = VD->getType();
175 if (Type->isLValueReferenceType())
176 Type = C.getPointerType(Type.getNonReferenceType());
177 else
178 Type = Type.getNonReferenceType();
179 SourceLocation Loc = VD->getLocation();
180 FieldDecl *Field;
181 if (SingleEscaped.count(VD)) {
182 Field = FieldDecl::Create(
183 C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,
184 C.getTrivialTypeSourceInfo(Type, SourceLocation()),
185 /*BW=*/nullptr, /*Mutable=*/false,
186 /*InitStyle=*/ICIS_NoInit);
187 Field->setAccess(AS_public);
188 if (VD->hasAttrs()) {
189 for (specific_attr_iterator<AlignedAttr> I(VD->getAttrs().begin()),
190 E(VD->getAttrs().end());
191 I != E; ++I)
192 Field->addAttr(*I);
193 }
194 } else {
195 llvm::APInt ArraySize(32, BufSize);
196 Type = C.getConstantArrayType(Type, ArraySize, nullptr, ArrayType::Normal,
197 0);
198 Field = FieldDecl::Create(
199 C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,
200 C.getTrivialTypeSourceInfo(Type, SourceLocation()),
201 /*BW=*/nullptr, /*Mutable=*/false,
202 /*InitStyle=*/ICIS_NoInit);
203 Field->setAccess(AS_public);
204 llvm::APInt Align(32, std::max(C.getDeclAlign(VD).getQuantity(),
205 static_cast<CharUnits::QuantityType>(
206 GlobalMemoryAlignment)));
207 Field->addAttr(AlignedAttr::CreateImplicit(
208 C, /*IsAlignmentExpr=*/true,
209 IntegerLiteral::Create(C, Align,
210 C.getIntTypeForBitwidth(32, /*Signed=*/0),
211 SourceLocation()),
212 {}, AttributeCommonInfo::AS_GNU, AlignedAttr::GNU_aligned));
213 }
214 GlobalizedRD->addDecl(Field);
215 MappedDeclsFields.try_emplace(VD, Field);
216 }
217 GlobalizedRD->completeDefinition();
218 return GlobalizedRD;
219}
220
221/// Get the list of variables that can escape their declaration context.
222class CheckVarsEscapingDeclContext final
223 : public ConstStmtVisitor<CheckVarsEscapingDeclContext> {
224 CodeGenFunction &CGF;
225 llvm::SetVector<const ValueDecl *> EscapedDecls;
226 llvm::SetVector<const ValueDecl *> EscapedVariableLengthDecls;
227 llvm::SmallPtrSet<const Decl *, 4> EscapedParameters;
228 RecordDecl *GlobalizedRD = nullptr;
229 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
230 bool AllEscaped = false;
231 bool IsForCombinedParallelRegion = false;
232
233 void markAsEscaped(const ValueDecl *VD) {
234 // Do not globalize declare target variables.
235 if (!isa<VarDecl>(VD) ||
236 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
237 return;
238 VD = cast<ValueDecl>(VD->getCanonicalDecl());
239 // Use user-specified allocation.
240 if (VD->hasAttrs() && VD->hasAttr<OMPAllocateDeclAttr>())
241 return;
242 // Variables captured by value must be globalized.
243 if (auto *CSI = CGF.CapturedStmtInfo) {
244 if (const FieldDecl *FD = CSI->lookup(cast<VarDecl>(VD))) {
245 // Check if need to capture the variable that was already captured by
246 // value in the outer region.
247 if (!IsForCombinedParallelRegion) {
248 if (!FD->hasAttrs())
249 return;
250 const auto *Attr = FD->getAttr<OMPCaptureKindAttr>();
251 if (!Attr)
252 return;
253 if (((Attr->getCaptureKind() != OMPC_map) &&
254 !isOpenMPPrivate(Attr->getCaptureKind())) ||
255 ((Attr->getCaptureKind() == OMPC_map) &&
256 !FD->getType()->isAnyPointerType()))
257 return;
258 }
259 if (!FD->getType()->isReferenceType()) {
260 assert(!VD->getType()->isVariablyModifiedType() &&((!VD->getType()->isVariablyModifiedType() && "Parameter captured by value with variably modified type"
) ? static_cast<void> (0) : __assert_fail ("!VD->getType()->isVariablyModifiedType() && \"Parameter captured by value with variably modified type\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 261, __PRETTY_FUNCTION__))
261 "Parameter captured by value with variably modified type")((!VD->getType()->isVariablyModifiedType() && "Parameter captured by value with variably modified type"
) ? static_cast<void> (0) : __assert_fail ("!VD->getType()->isVariablyModifiedType() && \"Parameter captured by value with variably modified type\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 261, __PRETTY_FUNCTION__))
;
262 EscapedParameters.insert(VD);
263 } else if (!IsForCombinedParallelRegion) {
264 return;
265 }
266 }
267 }
268 if ((!CGF.CapturedStmtInfo ||
269 (IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) &&
270 VD->getType()->isReferenceType())
271 // Do not globalize variables with reference type.
272 return;
273 if (VD->getType()->isVariablyModifiedType())
274 EscapedVariableLengthDecls.insert(VD);
275 else
276 EscapedDecls.insert(VD);
277 }
278
279 void VisitValueDecl(const ValueDecl *VD) {
280 if (VD->getType()->isLValueReferenceType())
281 markAsEscaped(VD);
282 if (const auto *VarD = dyn_cast<VarDecl>(VD)) {
283 if (!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {
284 const bool SavedAllEscaped = AllEscaped;
285 AllEscaped = VD->getType()->isLValueReferenceType();
286 Visit(VarD->getInit());
287 AllEscaped = SavedAllEscaped;
288 }
289 }
290 }
291 void VisitOpenMPCapturedStmt(const CapturedStmt *S,
292 ArrayRef<OMPClause *> Clauses,
293 bool IsCombinedParallelRegion) {
294 if (!S)
295 return;
296 for (const CapturedStmt::Capture &C : S->captures()) {
297 if (C.capturesVariable() && !C.capturesVariableByCopy()) {
298 const ValueDecl *VD = C.getCapturedVar();
299 bool SavedIsForCombinedParallelRegion = IsForCombinedParallelRegion;
300 if (IsCombinedParallelRegion) {
301 // Check if the variable is privatized in the combined construct and
302 // those private copies must be shared in the inner parallel
303 // directive.
304 IsForCombinedParallelRegion = false;
305 for (const OMPClause *C : Clauses) {
306 if (!isOpenMPPrivate(C->getClauseKind()) ||
307 C->getClauseKind() == OMPC_reduction ||
308 C->getClauseKind() == OMPC_linear ||
309 C->getClauseKind() == OMPC_private)
310 continue;
311 ArrayRef<const Expr *> Vars;
312 if (const auto *PC = dyn_cast<OMPFirstprivateClause>(C))
313 Vars = PC->getVarRefs();
314 else if (const auto *PC = dyn_cast<OMPLastprivateClause>(C))
315 Vars = PC->getVarRefs();
316 else
317 llvm_unreachable("Unexpected clause.")::llvm::llvm_unreachable_internal("Unexpected clause.", "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 317)
;
318 for (const auto *E : Vars) {
319 const Decl *D =
320 cast<DeclRefExpr>(E)->getDecl()->getCanonicalDecl();
321 if (D == VD->getCanonicalDecl()) {
322 IsForCombinedParallelRegion = true;
323 break;
324 }
325 }
326 if (IsForCombinedParallelRegion)
327 break;
328 }
329 }
330 markAsEscaped(VD);
331 if (isa<OMPCapturedExprDecl>(VD))
332 VisitValueDecl(VD);
333 IsForCombinedParallelRegion = SavedIsForCombinedParallelRegion;
334 }
335 }
336 }
337
338 void buildRecordForGlobalizedVars(bool IsInTTDRegion) {
339 assert(!GlobalizedRD &&((!GlobalizedRD && "Record for globalized variables is built already."
) ? static_cast<void> (0) : __assert_fail ("!GlobalizedRD && \"Record for globalized variables is built already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 340, __PRETTY_FUNCTION__))
340 "Record for globalized variables is built already.")((!GlobalizedRD && "Record for globalized variables is built already."
) ? static_cast<void> (0) : __assert_fail ("!GlobalizedRD && \"Record for globalized variables is built already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 340, __PRETTY_FUNCTION__))
;
341 ArrayRef<const ValueDecl *> EscapedDeclsForParallel, EscapedDeclsForTeams;
342 unsigned WarpSize = CGF.getTarget().getGridValue(llvm::omp::GV_Warp_Size);
343 if (IsInTTDRegion)
344 EscapedDeclsForTeams = EscapedDecls.getArrayRef();
345 else
346 EscapedDeclsForParallel = EscapedDecls.getArrayRef();
347 GlobalizedRD = ::buildRecordForGlobalizedVars(
348 CGF.getContext(), EscapedDeclsForParallel, EscapedDeclsForTeams,
349 MappedDeclsFields, WarpSize);
350 }
351
352public:
353 CheckVarsEscapingDeclContext(CodeGenFunction &CGF,
354 ArrayRef<const ValueDecl *> TeamsReductions)
355 : CGF(CGF), EscapedDecls(TeamsReductions.begin(), TeamsReductions.end()) {
356 }
357 virtual ~CheckVarsEscapingDeclContext() = default;
358 void VisitDeclStmt(const DeclStmt *S) {
359 if (!S)
360 return;
361 for (const Decl *D : S->decls())
362 if (const auto *VD = dyn_cast_or_null<ValueDecl>(D))
363 VisitValueDecl(VD);
364 }
365 void VisitOMPExecutableDirective(const OMPExecutableDirective *D) {
366 if (!D)
367 return;
368 if (!D->hasAssociatedStmt())
369 return;
370 if (const auto *S =
371 dyn_cast_or_null<CapturedStmt>(D->getAssociatedStmt())) {
372 // Do not analyze directives that do not actually require capturing,
373 // like `omp for` or `omp simd` directives.
374 llvm::SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
375 getOpenMPCaptureRegions(CaptureRegions, D->getDirectiveKind());
376 if (CaptureRegions.size() == 1 && CaptureRegions.back() == OMPD_unknown) {
377 VisitStmt(S->getCapturedStmt());
378 return;
379 }
380 VisitOpenMPCapturedStmt(
381 S, D->clauses(),
382 CaptureRegions.back() == OMPD_parallel &&
383 isOpenMPDistributeDirective(D->getDirectiveKind()));
384 }
385 }
386 void VisitCapturedStmt(const CapturedStmt *S) {
387 if (!S)
388 return;
389 for (const CapturedStmt::Capture &C : S->captures()) {
390 if (C.capturesVariable() && !C.capturesVariableByCopy()) {
391 const ValueDecl *VD = C.getCapturedVar();
392 markAsEscaped(VD);
393 if (isa<OMPCapturedExprDecl>(VD))
394 VisitValueDecl(VD);
395 }
396 }
397 }
398 void VisitLambdaExpr(const LambdaExpr *E) {
399 if (!E)
400 return;
401 for (const LambdaCapture &C : E->captures()) {
402 if (C.capturesVariable()) {
403 if (C.getCaptureKind() == LCK_ByRef) {
404 const ValueDecl *VD = C.getCapturedVar();
405 markAsEscaped(VD);
406 if (E->isInitCapture(&C) || isa<OMPCapturedExprDecl>(VD))
407 VisitValueDecl(VD);
408 }
409 }
410 }
411 }
412 void VisitBlockExpr(const BlockExpr *E) {
413 if (!E)
414 return;
415 for (const BlockDecl::Capture &C : E->getBlockDecl()->captures()) {
416 if (C.isByRef()) {
417 const VarDecl *VD = C.getVariable();
418 markAsEscaped(VD);
419 if (isa<OMPCapturedExprDecl>(VD) || VD->isInitCapture())
420 VisitValueDecl(VD);
421 }
422 }
423 }
424 void VisitCallExpr(const CallExpr *E) {
425 if (!E)
426 return;
427 for (const Expr *Arg : E->arguments()) {
428 if (!Arg)
429 continue;
430 if (Arg->isLValue()) {
431 const bool SavedAllEscaped = AllEscaped;
432 AllEscaped = true;
433 Visit(Arg);
434 AllEscaped = SavedAllEscaped;
435 } else {
436 Visit(Arg);
437 }
438 }
439 Visit(E->getCallee());
440 }
441 void VisitDeclRefExpr(const DeclRefExpr *E) {
442 if (!E)
443 return;
444 const ValueDecl *VD = E->getDecl();
445 if (AllEscaped)
446 markAsEscaped(VD);
447 if (isa<OMPCapturedExprDecl>(VD))
448 VisitValueDecl(VD);
449 else if (const auto *VarD = dyn_cast<VarDecl>(VD))
450 if (VarD->isInitCapture())
451 VisitValueDecl(VD);
452 }
453 void VisitUnaryOperator(const UnaryOperator *E) {
454 if (!E)
455 return;
456 if (E->getOpcode() == UO_AddrOf) {
457 const bool SavedAllEscaped = AllEscaped;
458 AllEscaped = true;
459 Visit(E->getSubExpr());
460 AllEscaped = SavedAllEscaped;
461 } else {
462 Visit(E->getSubExpr());
463 }
464 }
465 void VisitImplicitCastExpr(const ImplicitCastExpr *E) {
466 if (!E)
467 return;
468 if (E->getCastKind() == CK_ArrayToPointerDecay) {
469 const bool SavedAllEscaped = AllEscaped;
470 AllEscaped = true;
471 Visit(E->getSubExpr());
472 AllEscaped = SavedAllEscaped;
473 } else {
474 Visit(E->getSubExpr());
475 }
476 }
477 void VisitExpr(const Expr *E) {
478 if (!E)
479 return;
480 bool SavedAllEscaped = AllEscaped;
481 if (!E->isLValue())
482 AllEscaped = false;
483 for (const Stmt *Child : E->children())
484 if (Child)
485 Visit(Child);
486 AllEscaped = SavedAllEscaped;
487 }
488 void VisitStmt(const Stmt *S) {
489 if (!S)
490 return;
491 for (const Stmt *Child : S->children())
492 if (Child)
493 Visit(Child);
494 }
495
496 /// Returns the record that handles all the escaped local variables and used
497 /// instead of their original storage.
498 const RecordDecl *getGlobalizedRecord(bool IsInTTDRegion) {
499 if (!GlobalizedRD)
500 buildRecordForGlobalizedVars(IsInTTDRegion);
501 return GlobalizedRD;
502 }
503
504 /// Returns the field in the globalized record for the escaped variable.
505 const FieldDecl *getFieldForGlobalizedVar(const ValueDecl *VD) const {
506 assert(GlobalizedRD &&((GlobalizedRD && "Record for globalized variables must be generated already."
) ? static_cast<void> (0) : __assert_fail ("GlobalizedRD && \"Record for globalized variables must be generated already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 507, __PRETTY_FUNCTION__))
507 "Record for globalized variables must be generated already.")((GlobalizedRD && "Record for globalized variables must be generated already."
) ? static_cast<void> (0) : __assert_fail ("GlobalizedRD && \"Record for globalized variables must be generated already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 507, __PRETTY_FUNCTION__))
;
508 auto I = MappedDeclsFields.find(VD);
509 if (I == MappedDeclsFields.end())
510 return nullptr;
511 return I->getSecond();
512 }
513
514 /// Returns the list of the escaped local variables/parameters.
515 ArrayRef<const ValueDecl *> getEscapedDecls() const {
516 return EscapedDecls.getArrayRef();
517 }
518
519 /// Checks if the escaped local variable is actually a parameter passed by
520 /// value.
521 const llvm::SmallPtrSetImpl<const Decl *> &getEscapedParameters() const {
522 return EscapedParameters;
523 }
524
525 /// Returns the list of the escaped variables with the variably modified
526 /// types.
527 ArrayRef<const ValueDecl *> getEscapedVariableLengthDecls() const {
528 return EscapedVariableLengthDecls.getArrayRef();
529 }
530};
531} // anonymous namespace
532
533/// Get the id of the warp in the block.
534/// We assume that the warp size is 32, which is always the case
535/// on the NVPTX device, to generate more efficient code.
536static llvm::Value *getNVPTXWarpID(CodeGenFunction &CGF) {
537 CGBuilderTy &Bld = CGF.Builder;
538 unsigned LaneIDBits =
539 CGF.getTarget().getGridValue(llvm::omp::GV_Warp_Size_Log2);
540 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
541 return Bld.CreateAShr(RT.getGPUThreadID(CGF), LaneIDBits, "nvptx_warp_id");
542}
543
544/// Get the id of the current lane in the Warp.
545/// We assume that the warp size is 32, which is always the case
546/// on the NVPTX device, to generate more efficient code.
547static llvm::Value *getNVPTXLaneID(CodeGenFunction &CGF) {
548 CGBuilderTy &Bld = CGF.Builder;
549 unsigned LaneIDMask = CGF.getContext().getTargetInfo().getGridValue(
550 llvm::omp::GV_Warp_Size_Log2_Mask);
551 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
552 return Bld.CreateAnd(RT.getGPUThreadID(CGF), Bld.getInt32(LaneIDMask),
553 "nvptx_lane_id");
554}
555
556/// Get the value of the thread_limit clause in the teams directive.
557/// For the 'generic' execution mode, the runtime encodes thread_limit in
558/// the launch parameters, always starting thread_limit+warpSize threads per
559/// CTA. The threads in the last warp are reserved for master execution.
560/// For the 'spmd' execution mode, all threads in a CTA are part of the team.
561static llvm::Value *getThreadLimit(CodeGenFunction &CGF,
562 bool IsInSPMDExecutionMode = false) {
563 CGBuilderTy &Bld = CGF.Builder;
564 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
565 return IsInSPMDExecutionMode
566 ? RT.getGPUNumThreads(CGF)
567 : Bld.CreateNUWSub(RT.getGPUNumThreads(CGF),
568 RT.getGPUWarpSize(CGF), "thread_limit");
569}
570
571/// Get the thread id of the OMP master thread.
572/// The master thread id is the first thread (lane) of the last warp in the
573/// GPU block. Warp size is assumed to be some power of 2.
574/// Thread id is 0 indexed.
575/// E.g: If NumThreads is 33, master id is 32.
576/// If NumThreads is 64, master id is 32.
577/// If NumThreads is 1024, master id is 992.
578static llvm::Value *getMasterThreadID(CodeGenFunction &CGF) {
579 CGBuilderTy &Bld = CGF.Builder;
580 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
581 llvm::Value *NumThreads = RT.getGPUNumThreads(CGF);
582 // We assume that the warp size is a power of 2.
583 llvm::Value *Mask = Bld.CreateNUWSub(RT.getGPUWarpSize(CGF), Bld.getInt32(1));
584
585 return Bld.CreateAnd(Bld.CreateNUWSub(NumThreads, Bld.getInt32(1)),
586 Bld.CreateNot(Mask), "master_tid");
587}
588
589CGOpenMPRuntimeGPU::WorkerFunctionState::WorkerFunctionState(
590 CodeGenModule &CGM, SourceLocation Loc)
591 : WorkerFn(nullptr), CGFI(CGM.getTypes().arrangeNullaryFunction()),
592 Loc(Loc) {
593 createWorkerFunction(CGM);
594}
595
596void CGOpenMPRuntimeGPU::WorkerFunctionState::createWorkerFunction(
597 CodeGenModule &CGM) {
598 // Create an worker function with no arguments.
599
600 WorkerFn = llvm::Function::Create(
601 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
602 /*placeholder=*/"_worker", &CGM.getModule());
603 CGM.SetInternalFunctionAttributes(GlobalDecl(), WorkerFn, CGFI);
604 WorkerFn->setDoesNotRecurse();
605}
606
607CGOpenMPRuntimeGPU::ExecutionMode
608CGOpenMPRuntimeGPU::getExecutionMode() const {
609 return CurrentExecutionMode;
610}
611
612static CGOpenMPRuntimeGPU::DataSharingMode
613getDataSharingMode(CodeGenModule &CGM) {
614 return CGM.getLangOpts().OpenMPCUDAMode ? CGOpenMPRuntimeGPU::CUDA
615 : CGOpenMPRuntimeGPU::Generic;
616}
617
618/// Check for inner (nested) SPMD construct, if any
619static bool hasNestedSPMDDirective(ASTContext &Ctx,
620 const OMPExecutableDirective &D) {
621 const auto *CS = D.getInnermostCapturedStmt();
622 const auto *Body =
623 CS->getCapturedStmt()->IgnoreContainers(/*IgnoreCaptured=*/true);
624 const Stmt *ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
625
626 if (const auto *NestedDir =
627 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
628 OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind();
629 switch (D.getDirectiveKind()) {
630 case OMPD_target:
631 if (isOpenMPParallelDirective(DKind))
632 return true;
633 if (DKind == OMPD_teams) {
634 Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
635 /*IgnoreCaptured=*/true);
636 if (!Body)
637 return false;
638 ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
639 if (const auto *NND =
640 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
641 DKind = NND->getDirectiveKind();
642 if (isOpenMPParallelDirective(DKind))
643 return true;
644 }
645 }
646 return false;
647 case OMPD_target_teams:
648 return isOpenMPParallelDirective(DKind);
649 case OMPD_target_simd:
650 case OMPD_target_parallel:
651 case OMPD_target_parallel_for:
652 case OMPD_target_parallel_for_simd:
653 case OMPD_target_teams_distribute:
654 case OMPD_target_teams_distribute_simd:
655 case OMPD_target_teams_distribute_parallel_for:
656 case OMPD_target_teams_distribute_parallel_for_simd:
657 case OMPD_parallel:
658 case OMPD_for:
659 case OMPD_parallel_for:
660 case OMPD_parallel_master:
661 case OMPD_parallel_sections:
662 case OMPD_for_simd:
663 case OMPD_parallel_for_simd:
664 case OMPD_cancel:
665 case OMPD_cancellation_point:
666 case OMPD_ordered:
667 case OMPD_threadprivate:
668 case OMPD_allocate:
669 case OMPD_task:
670 case OMPD_simd:
671 case OMPD_sections:
672 case OMPD_section:
673 case OMPD_single:
674 case OMPD_master:
675 case OMPD_critical:
676 case OMPD_taskyield:
677 case OMPD_barrier:
678 case OMPD_taskwait:
679 case OMPD_taskgroup:
680 case OMPD_atomic:
681 case OMPD_flush:
682 case OMPD_depobj:
683 case OMPD_scan:
684 case OMPD_teams:
685 case OMPD_target_data:
686 case OMPD_target_exit_data:
687 case OMPD_target_enter_data:
688 case OMPD_distribute:
689 case OMPD_distribute_simd:
690 case OMPD_distribute_parallel_for:
691 case OMPD_distribute_parallel_for_simd:
692 case OMPD_teams_distribute:
693 case OMPD_teams_distribute_simd:
694 case OMPD_teams_distribute_parallel_for:
695 case OMPD_teams_distribute_parallel_for_simd:
696 case OMPD_target_update:
697 case OMPD_declare_simd:
698 case OMPD_declare_variant:
699 case OMPD_begin_declare_variant:
700 case OMPD_end_declare_variant:
701 case OMPD_declare_target:
702 case OMPD_end_declare_target:
703 case OMPD_declare_reduction:
704 case OMPD_declare_mapper:
705 case OMPD_taskloop:
706 case OMPD_taskloop_simd:
707 case OMPD_master_taskloop:
708 case OMPD_master_taskloop_simd:
709 case OMPD_parallel_master_taskloop:
710 case OMPD_parallel_master_taskloop_simd:
711 case OMPD_requires:
712 case OMPD_unknown:
713 default:
714 llvm_unreachable("Unexpected directive.")::llvm::llvm_unreachable_internal("Unexpected directive.", "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 714)
;
715 }
716 }
717
718 return false;
719}
720
721static bool supportsSPMDExecutionMode(ASTContext &Ctx,
722 const OMPExecutableDirective &D) {
723 OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
724 switch (DirectiveKind) {
725 case OMPD_target:
726 case OMPD_target_teams:
727 return hasNestedSPMDDirective(Ctx, D);
728 case OMPD_target_parallel:
729 case OMPD_target_parallel_for:
730 case OMPD_target_parallel_for_simd:
731 case OMPD_target_teams_distribute_parallel_for:
732 case OMPD_target_teams_distribute_parallel_for_simd:
733 case OMPD_target_simd:
734 case OMPD_target_teams_distribute_simd:
735 return true;
736 case OMPD_target_teams_distribute:
737 return false;
738 case OMPD_parallel:
739 case OMPD_for:
740 case OMPD_parallel_for:
741 case OMPD_parallel_master:
742 case OMPD_parallel_sections:
743 case OMPD_for_simd:
744 case OMPD_parallel_for_simd:
745 case OMPD_cancel:
746 case OMPD_cancellation_point:
747 case OMPD_ordered:
748 case OMPD_threadprivate:
749 case OMPD_allocate:
750 case OMPD_task:
751 case OMPD_simd:
752 case OMPD_sections:
753 case OMPD_section:
754 case OMPD_single:
755 case OMPD_master:
756 case OMPD_critical:
757 case OMPD_taskyield:
758 case OMPD_barrier:
759 case OMPD_taskwait:
760 case OMPD_taskgroup:
761 case OMPD_atomic:
762 case OMPD_flush:
763 case OMPD_depobj:
764 case OMPD_scan:
765 case OMPD_teams:
766 case OMPD_target_data:
767 case OMPD_target_exit_data:
768 case OMPD_target_enter_data:
769 case OMPD_distribute:
770 case OMPD_distribute_simd:
771 case OMPD_distribute_parallel_for:
772 case OMPD_distribute_parallel_for_simd:
773 case OMPD_teams_distribute:
774 case OMPD_teams_distribute_simd:
775 case OMPD_teams_distribute_parallel_for:
776 case OMPD_teams_distribute_parallel_for_simd:
777 case OMPD_target_update:
778 case OMPD_declare_simd:
779 case OMPD_declare_variant:
780 case OMPD_begin_declare_variant:
781 case OMPD_end_declare_variant:
782 case OMPD_declare_target:
783 case OMPD_end_declare_target:
784 case OMPD_declare_reduction:
785 case OMPD_declare_mapper:
786 case OMPD_taskloop:
787 case OMPD_taskloop_simd:
788 case OMPD_master_taskloop:
789 case OMPD_master_taskloop_simd:
790 case OMPD_parallel_master_taskloop:
791 case OMPD_parallel_master_taskloop_simd:
792 case OMPD_requires:
793 case OMPD_unknown:
794 default:
795 break;
796 }
797 llvm_unreachable(::llvm::llvm_unreachable_internal("Unknown programming model for OpenMP directive on NVPTX target."
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 798)
798 "Unknown programming model for OpenMP directive on NVPTX target.")::llvm::llvm_unreachable_internal("Unknown programming model for OpenMP directive on NVPTX target."
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 798)
;
799}
800
801/// Check if the directive is loops based and has schedule clause at all or has
802/// static scheduling.
803static bool hasStaticScheduling(const OMPExecutableDirective &D) {
804 assert(isOpenMPWorksharingDirective(D.getDirectiveKind()) &&((isOpenMPWorksharingDirective(D.getDirectiveKind()) &&
isOpenMPLoopDirective(D.getDirectiveKind()) && "Expected loop-based directive."
) ? static_cast<void> (0) : __assert_fail ("isOpenMPWorksharingDirective(D.getDirectiveKind()) && isOpenMPLoopDirective(D.getDirectiveKind()) && \"Expected loop-based directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 806, __PRETTY_FUNCTION__))
805 isOpenMPLoopDirective(D.getDirectiveKind()) &&((isOpenMPWorksharingDirective(D.getDirectiveKind()) &&
isOpenMPLoopDirective(D.getDirectiveKind()) && "Expected loop-based directive."
) ? static_cast<void> (0) : __assert_fail ("isOpenMPWorksharingDirective(D.getDirectiveKind()) && isOpenMPLoopDirective(D.getDirectiveKind()) && \"Expected loop-based directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 806, __PRETTY_FUNCTION__))
806 "Expected loop-based directive.")((isOpenMPWorksharingDirective(D.getDirectiveKind()) &&
isOpenMPLoopDirective(D.getDirectiveKind()) && "Expected loop-based directive."
) ? static_cast<void> (0) : __assert_fail ("isOpenMPWorksharingDirective(D.getDirectiveKind()) && isOpenMPLoopDirective(D.getDirectiveKind()) && \"Expected loop-based directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 806, __PRETTY_FUNCTION__))
;
807 return !D.hasClausesOfKind<OMPOrderedClause>() &&
808 (!D.hasClausesOfKind<OMPScheduleClause>() ||
809 llvm::any_of(D.getClausesOfKind<OMPScheduleClause>(),
810 [](const OMPScheduleClause *C) {
811 return C->getScheduleKind() == OMPC_SCHEDULE_static;
812 }));
813}
814
815/// Check for inner (nested) lightweight runtime construct, if any
816static bool hasNestedLightweightDirective(ASTContext &Ctx,
817 const OMPExecutableDirective &D) {
818 assert(supportsSPMDExecutionMode(Ctx, D) && "Expected SPMD mode directive.")((supportsSPMDExecutionMode(Ctx, D) && "Expected SPMD mode directive."
) ? static_cast<void> (0) : __assert_fail ("supportsSPMDExecutionMode(Ctx, D) && \"Expected SPMD mode directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 818, __PRETTY_FUNCTION__))
;
819 const auto *CS = D.getInnermostCapturedStmt();
820 const auto *Body =
821 CS->getCapturedStmt()->IgnoreContainers(/*IgnoreCaptured=*/true);
822 const Stmt *ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
823
824 if (const auto *NestedDir =
825 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
826 OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind();
827 switch (D.getDirectiveKind()) {
828 case OMPD_target:
829 if (isOpenMPParallelDirective(DKind) &&
830 isOpenMPWorksharingDirective(DKind) && isOpenMPLoopDirective(DKind) &&
831 hasStaticScheduling(*NestedDir))
832 return true;
833 if (DKind == OMPD_teams_distribute_simd || DKind == OMPD_simd)
834 return true;
835 if (DKind == OMPD_parallel) {
836 Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
837 /*IgnoreCaptured=*/true);
838 if (!Body)
839 return false;
840 ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
841 if (const auto *NND =
842 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
843 DKind = NND->getDirectiveKind();
844 if (isOpenMPWorksharingDirective(DKind) &&
845 isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
846 return true;
847 }
848 } else if (DKind == OMPD_teams) {
849 Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
850 /*IgnoreCaptured=*/true);
851 if (!Body)
852 return false;
853 ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
854 if (const auto *NND =
855 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
856 DKind = NND->getDirectiveKind();
857 if (isOpenMPParallelDirective(DKind) &&
858 isOpenMPWorksharingDirective(DKind) &&
859 isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
860 return true;
861 if (DKind == OMPD_parallel) {
862 Body = NND->getInnermostCapturedStmt()->IgnoreContainers(
863 /*IgnoreCaptured=*/true);
864 if (!Body)
865 return false;
866 ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
867 if (const auto *NND =
868 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
869 DKind = NND->getDirectiveKind();
870 if (isOpenMPWorksharingDirective(DKind) &&
871 isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
872 return true;
873 }
874 }
875 }
876 }
877 return false;
878 case OMPD_target_teams:
879 if (isOpenMPParallelDirective(DKind) &&
880 isOpenMPWorksharingDirective(DKind) && isOpenMPLoopDirective(DKind) &&
881 hasStaticScheduling(*NestedDir))
882 return true;
883 if (DKind == OMPD_distribute_simd || DKind == OMPD_simd)
884 return true;
885 if (DKind == OMPD_parallel) {
886 Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
887 /*IgnoreCaptured=*/true);
888 if (!Body)
889 return false;
890 ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
891 if (const auto *NND =
892 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
893 DKind = NND->getDirectiveKind();
894 if (isOpenMPWorksharingDirective(DKind) &&
895 isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
896 return true;
897 }
898 }
899 return false;
900 case OMPD_target_parallel:
901 if (DKind == OMPD_simd)
902 return true;
903 return isOpenMPWorksharingDirective(DKind) &&
904 isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NestedDir);
905 case OMPD_target_teams_distribute:
906 case OMPD_target_simd:
907 case OMPD_target_parallel_for:
908 case OMPD_target_parallel_for_simd:
909 case OMPD_target_teams_distribute_simd:
910 case OMPD_target_teams_distribute_parallel_for:
911 case OMPD_target_teams_distribute_parallel_for_simd:
912 case OMPD_parallel:
913 case OMPD_for:
914 case OMPD_parallel_for:
915 case OMPD_parallel_master:
916 case OMPD_parallel_sections:
917 case OMPD_for_simd:
918 case OMPD_parallel_for_simd:
919 case OMPD_cancel:
920 case OMPD_cancellation_point:
921 case OMPD_ordered:
922 case OMPD_threadprivate:
923 case OMPD_allocate:
924 case OMPD_task:
925 case OMPD_simd:
926 case OMPD_sections:
927 case OMPD_section:
928 case OMPD_single:
929 case OMPD_master:
930 case OMPD_critical:
931 case OMPD_taskyield:
932 case OMPD_barrier:
933 case OMPD_taskwait:
934 case OMPD_taskgroup:
935 case OMPD_atomic:
936 case OMPD_flush:
937 case OMPD_depobj:
938 case OMPD_scan:
939 case OMPD_teams:
940 case OMPD_target_data:
941 case OMPD_target_exit_data:
942 case OMPD_target_enter_data:
943 case OMPD_distribute:
944 case OMPD_distribute_simd:
945 case OMPD_distribute_parallel_for:
946 case OMPD_distribute_parallel_for_simd:
947 case OMPD_teams_distribute:
948 case OMPD_teams_distribute_simd:
949 case OMPD_teams_distribute_parallel_for:
950 case OMPD_teams_distribute_parallel_for_simd:
951 case OMPD_target_update:
952 case OMPD_declare_simd:
953 case OMPD_declare_variant:
954 case OMPD_begin_declare_variant:
955 case OMPD_end_declare_variant:
956 case OMPD_declare_target:
957 case OMPD_end_declare_target:
958 case OMPD_declare_reduction:
959 case OMPD_declare_mapper:
960 case OMPD_taskloop:
961 case OMPD_taskloop_simd:
962 case OMPD_master_taskloop:
963 case OMPD_master_taskloop_simd:
964 case OMPD_parallel_master_taskloop:
965 case OMPD_parallel_master_taskloop_simd:
966 case OMPD_requires:
967 case OMPD_unknown:
968 default:
969 llvm_unreachable("Unexpected directive.")::llvm::llvm_unreachable_internal("Unexpected directive.", "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 969)
;
970 }
971 }
972
973 return false;
974}
975
976/// Checks if the construct supports lightweight runtime. It must be SPMD
977/// construct + inner loop-based construct with static scheduling.
978static bool supportsLightweightRuntime(ASTContext &Ctx,
979 const OMPExecutableDirective &D) {
980 if (!supportsSPMDExecutionMode(Ctx, D))
981 return false;
982 OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
983 switch (DirectiveKind) {
984 case OMPD_target:
985 case OMPD_target_teams:
986 case OMPD_target_parallel:
987 return hasNestedLightweightDirective(Ctx, D);
988 case OMPD_target_parallel_for:
989 case OMPD_target_parallel_for_simd:
990 case OMPD_target_teams_distribute_parallel_for:
991 case OMPD_target_teams_distribute_parallel_for_simd:
992 // (Last|First)-privates must be shared in parallel region.
993 return hasStaticScheduling(D);
994 case OMPD_target_simd:
995 case OMPD_target_teams_distribute_simd:
996 return true;
997 case OMPD_target_teams_distribute:
998 return false;
999 case OMPD_parallel:
1000 case OMPD_for:
1001 case OMPD_parallel_for:
1002 case OMPD_parallel_master:
1003 case OMPD_parallel_sections:
1004 case OMPD_for_simd:
1005 case OMPD_parallel_for_simd:
1006 case OMPD_cancel:
1007 case OMPD_cancellation_point:
1008 case OMPD_ordered:
1009 case OMPD_threadprivate:
1010 case OMPD_allocate:
1011 case OMPD_task:
1012 case OMPD_simd:
1013 case OMPD_sections:
1014 case OMPD_section:
1015 case OMPD_single:
1016 case OMPD_master:
1017 case OMPD_critical:
1018 case OMPD_taskyield:
1019 case OMPD_barrier:
1020 case OMPD_taskwait:
1021 case OMPD_taskgroup:
1022 case OMPD_atomic:
1023 case OMPD_flush:
1024 case OMPD_depobj:
1025 case OMPD_scan:
1026 case OMPD_teams:
1027 case OMPD_target_data:
1028 case OMPD_target_exit_data:
1029 case OMPD_target_enter_data:
1030 case OMPD_distribute:
1031 case OMPD_distribute_simd:
1032 case OMPD_distribute_parallel_for:
1033 case OMPD_distribute_parallel_for_simd:
1034 case OMPD_teams_distribute:
1035 case OMPD_teams_distribute_simd:
1036 case OMPD_teams_distribute_parallel_for:
1037 case OMPD_teams_distribute_parallel_for_simd:
1038 case OMPD_target_update:
1039 case OMPD_declare_simd:
1040 case OMPD_declare_variant:
1041 case OMPD_begin_declare_variant:
1042 case OMPD_end_declare_variant:
1043 case OMPD_declare_target:
1044 case OMPD_end_declare_target:
1045 case OMPD_declare_reduction:
1046 case OMPD_declare_mapper:
1047 case OMPD_taskloop:
1048 case OMPD_taskloop_simd:
1049 case OMPD_master_taskloop:
1050 case OMPD_master_taskloop_simd:
1051 case OMPD_parallel_master_taskloop:
1052 case OMPD_parallel_master_taskloop_simd:
1053 case OMPD_requires:
1054 case OMPD_unknown:
1055 default:
1056 break;
1057 }
1058 llvm_unreachable(::llvm::llvm_unreachable_internal("Unknown programming model for OpenMP directive on NVPTX target."
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1059)
1059 "Unknown programming model for OpenMP directive on NVPTX target.")::llvm::llvm_unreachable_internal("Unknown programming model for OpenMP directive on NVPTX target."
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1059)
;
1060}
1061
1062void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
1063 StringRef ParentName,
1064 llvm::Function *&OutlinedFn,
1065 llvm::Constant *&OutlinedFnID,
1066 bool IsOffloadEntry,
1067 const RegionCodeGenTy &CodeGen) {
1068 ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode);
1069 EntryFunctionState EST;
1070 WorkerFunctionState WST(CGM, D.getBeginLoc());
1071 Work.clear();
1072 WrapperFunctionsMap.clear();
1073
1074 // Emit target region as a standalone region.
1075 class NVPTXPrePostActionTy : public PrePostActionTy {
1076 CGOpenMPRuntimeGPU::EntryFunctionState &EST;
1077 CGOpenMPRuntimeGPU::WorkerFunctionState &WST;
1078
1079 public:
1080 NVPTXPrePostActionTy(CGOpenMPRuntimeGPU::EntryFunctionState &EST,
1081 CGOpenMPRuntimeGPU::WorkerFunctionState &WST)
1082 : EST(EST), WST(WST) {}
1083 void Enter(CodeGenFunction &CGF) override {
1084 auto &RT =
1085 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1086 RT.emitNonSPMDEntryHeader(CGF, EST, WST);
1087 // Skip target region initialization.
1088 RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
1089 }
1090 void Exit(CodeGenFunction &CGF) override {
1091 auto &RT =
1092 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1093 RT.clearLocThreadIdInsertPt(CGF);
1094 RT.emitNonSPMDEntryFooter(CGF, EST);
1095 }
1096 } Action(EST, WST);
1097 CodeGen.setAction(Action);
1098 IsInTTDRegion = true;
1099 // Reserve place for the globalized memory.
1100 GlobalizedRecords.emplace_back();
1101 if (!KernelStaticGlobalized) {
1102 KernelStaticGlobalized = new llvm::GlobalVariable(
1103 CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
1104 llvm::GlobalValue::InternalLinkage,
1105 llvm::UndefValue::get(CGM.VoidPtrTy),
1106 "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
1107 llvm::GlobalValue::NotThreadLocal,
1108 CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
1109 }
1110 emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
1111 IsOffloadEntry, CodeGen);
1112 IsInTTDRegion = false;
1113
1114 // Now change the name of the worker function to correspond to this target
1115 // region's entry function.
1116 WST.WorkerFn->setName(Twine(OutlinedFn->getName(), "_worker"));
1117
1118 // Create the worker function
1119 emitWorkerFunction(WST);
1120}
1121
1122// Setup NVPTX threads for master-worker OpenMP scheme.
1123void CGOpenMPRuntimeGPU::emitNonSPMDEntryHeader(CodeGenFunction &CGF,
1124 EntryFunctionState &EST,
1125 WorkerFunctionState &WST) {
1126 CGBuilderTy &Bld = CGF.Builder;
1127
1128 llvm::BasicBlock *WorkerBB = CGF.createBasicBlock(".worker");
1129 llvm::BasicBlock *MasterCheckBB = CGF.createBasicBlock(".mastercheck");
1130 llvm::BasicBlock *MasterBB = CGF.createBasicBlock(".master");
1131 EST.ExitBB = CGF.createBasicBlock(".exit");
1132
1133 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1134 llvm::Value *IsWorker =
1135 Bld.CreateICmpULT(RT.getGPUThreadID(CGF), getThreadLimit(CGF));
1136 Bld.CreateCondBr(IsWorker, WorkerBB, MasterCheckBB);
1137
1138 CGF.EmitBlock(WorkerBB);
1139 emitCall(CGF, WST.Loc, WST.WorkerFn);
1140 CGF.EmitBranch(EST.ExitBB);
1141
1142 CGF.EmitBlock(MasterCheckBB);
1143 llvm::Value *IsMaster =
1144 Bld.CreateICmpEQ(RT.getGPUThreadID(CGF), getMasterThreadID(CGF));
1145 Bld.CreateCondBr(IsMaster, MasterBB, EST.ExitBB);
1146
1147 CGF.EmitBlock(MasterBB);
1148 IsInTargetMasterThreadRegion = true;
1149 // SEQUENTIAL (MASTER) REGION START
1150 // First action in sequential region:
1151 // Initialize the state of the OpenMP runtime library on the GPU.
1152 // TODO: Optimize runtime initialization and pass in correct value.
1153 llvm::Value *Args[] = {getThreadLimit(CGF),
1154 Bld.getInt16(/*RequiresOMPRuntime=*/1)};
1155 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1156 CGM.getModule(), OMPRTL___kmpc_kernel_init),
1157 Args);
1158
1159 // For data sharing, we need to initialize the stack.
1160 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1161 CGM.getModule(), OMPRTL___kmpc_data_sharing_init_stack));
1162
1163 emitGenericVarsProlog(CGF, WST.Loc);
1164}
1165
1166void CGOpenMPRuntimeGPU::emitNonSPMDEntryFooter(CodeGenFunction &CGF,
1167 EntryFunctionState &EST) {
1168 IsInTargetMasterThreadRegion = false;
1169 if (!CGF.HaveInsertPoint())
1170 return;
1171
1172 emitGenericVarsEpilog(CGF);
1173
1174 if (!EST.ExitBB)
1175 EST.ExitBB = CGF.createBasicBlock(".exit");
1176
1177 llvm::BasicBlock *TerminateBB = CGF.createBasicBlock(".termination.notifier");
1178 CGF.EmitBranch(TerminateBB);
1179
1180 CGF.EmitBlock(TerminateBB);
1181 // Signal termination condition.
1182 // TODO: Optimize runtime initialization and pass in correct value.
1183 llvm::Value *Args[] = {CGF.Builder.getInt16(/*IsOMPRuntimeInitialized=*/1)};
1184 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1185 CGM.getModule(), OMPRTL___kmpc_kernel_deinit),
1186 Args);
1187 // Barrier to terminate worker threads.
1188 syncCTAThreads(CGF);
1189 // Master thread jumps to exit point.
1190 CGF.EmitBranch(EST.ExitBB);
1191
1192 CGF.EmitBlock(EST.ExitBB);
1193 EST.ExitBB = nullptr;
1194}
1195
1196void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
1197 StringRef ParentName,
1198 llvm::Function *&OutlinedFn,
1199 llvm::Constant *&OutlinedFnID,
1200 bool IsOffloadEntry,
1201 const RegionCodeGenTy &CodeGen) {
1202 ExecutionRuntimeModesRAII ModeRAII(
1203 CurrentExecutionMode, RequiresFullRuntime,
1204 CGM.getLangOpts().OpenMPCUDAForceFullRuntime ||
1205 !supportsLightweightRuntime(CGM.getContext(), D));
1206 EntryFunctionState EST;
1207
1208 // Emit target region as a standalone region.
1209 class NVPTXPrePostActionTy : public PrePostActionTy {
1210 CGOpenMPRuntimeGPU &RT;
1211 CGOpenMPRuntimeGPU::EntryFunctionState &EST;
1212 const OMPExecutableDirective &D;
1213
1214 public:
1215 NVPTXPrePostActionTy(CGOpenMPRuntimeGPU &RT,
1216 CGOpenMPRuntimeGPU::EntryFunctionState &EST,
1217 const OMPExecutableDirective &D)
1218 : RT(RT), EST(EST), D(D) {}
1219 void Enter(CodeGenFunction &CGF) override {
1220 RT.emitSPMDEntryHeader(CGF, EST, D);
1221 // Skip target region initialization.
1222 RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
1223 }
1224 void Exit(CodeGenFunction &CGF) override {
1225 RT.clearLocThreadIdInsertPt(CGF);
1226 RT.emitSPMDEntryFooter(CGF, EST);
1227 }
1228 } Action(*this, EST, D);
1229 CodeGen.setAction(Action);
1230 IsInTTDRegion = true;
1231 // Reserve place for the globalized memory.
1232 GlobalizedRecords.emplace_back();
1233 if (!KernelStaticGlobalized) {
1234 KernelStaticGlobalized = new llvm::GlobalVariable(
1235 CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
1236 llvm::GlobalValue::InternalLinkage,
1237 llvm::UndefValue::get(CGM.VoidPtrTy),
1238 "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
1239 llvm::GlobalValue::NotThreadLocal,
1240 CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
1241 }
1242 emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
1243 IsOffloadEntry, CodeGen);
1244 IsInTTDRegion = false;
1245}
1246
1247void CGOpenMPRuntimeGPU::emitSPMDEntryHeader(
1248 CodeGenFunction &CGF, EntryFunctionState &EST,
1249 const OMPExecutableDirective &D) {
1250 CGBuilderTy &Bld = CGF.Builder;
1251
1252 // Setup BBs in entry function.
1253 llvm::BasicBlock *ExecuteBB = CGF.createBasicBlock(".execute");
1254 EST.ExitBB = CGF.createBasicBlock(".exit");
1255
1256 llvm::Value *Args[] = {getThreadLimit(CGF, /*IsInSPMDExecutionMode=*/true),
1257 /*RequiresOMPRuntime=*/
1258 Bld.getInt16(RequiresFullRuntime ? 1 : 0)};
1259 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1260 CGM.getModule(), OMPRTL___kmpc_spmd_kernel_init),
1261 Args);
1262
1263 if (RequiresFullRuntime) {
1264 // For data sharing, we need to initialize the stack.
1265 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1266 CGM.getModule(), OMPRTL___kmpc_data_sharing_init_stack_spmd));
1267 }
1268
1269 CGF.EmitBranch(ExecuteBB);
1270
1271 CGF.EmitBlock(ExecuteBB);
1272
1273 IsInTargetMasterThreadRegion = true;
1274}
1275
1276void CGOpenMPRuntimeGPU::emitSPMDEntryFooter(CodeGenFunction &CGF,
1277 EntryFunctionState &EST) {
1278 IsInTargetMasterThreadRegion = false;
1279 if (!CGF.HaveInsertPoint())
1280 return;
1281
1282 if (!EST.ExitBB)
1283 EST.ExitBB = CGF.createBasicBlock(".exit");
1284
1285 llvm::BasicBlock *OMPDeInitBB = CGF.createBasicBlock(".omp.deinit");
1286 CGF.EmitBranch(OMPDeInitBB);
1287
1288 CGF.EmitBlock(OMPDeInitBB);
1289 // DeInitialize the OMP state in the runtime; called by all active threads.
1290 llvm::Value *Args[] = {/*RequiresOMPRuntime=*/
1291 CGF.Builder.getInt16(RequiresFullRuntime ? 1 : 0)};
1292 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1293 CGM.getModule(), OMPRTL___kmpc_spmd_kernel_deinit_v2),
1294 Args);
1295 CGF.EmitBranch(EST.ExitBB);
1296
1297 CGF.EmitBlock(EST.ExitBB);
1298 EST.ExitBB = nullptr;
1299}
1300
1301// Create a unique global variable to indicate the execution mode of this target
1302// region. The execution mode is either 'generic', or 'spmd' depending on the
1303// target directive. This variable is picked up by the offload library to setup
1304// the device appropriately before kernel launch. If the execution mode is
1305// 'generic', the runtime reserves one warp for the master, otherwise, all
1306// warps participate in parallel work.
1307static void setPropertyExecutionMode(CodeGenModule &CGM, StringRef Name,
1308 bool Mode) {
1309 auto *GVMode =
1310 new llvm::GlobalVariable(CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
1311 llvm::GlobalValue::WeakAnyLinkage,
1312 llvm::ConstantInt::get(CGM.Int8Ty, Mode ? 0 : 1),
1313 Twine(Name, "_exec_mode"));
1314 CGM.addCompilerUsedGlobal(GVMode);
1315}
1316
1317void CGOpenMPRuntimeGPU::emitWorkerFunction(WorkerFunctionState &WST) {
1318 ASTContext &Ctx = CGM.getContext();
1319
1320 CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
1321 CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, WST.WorkerFn, WST.CGFI, {},
1322 WST.Loc, WST.Loc);
1323 emitWorkerLoop(CGF, WST);
1324 CGF.FinishFunction();
1325}
1326
1327void CGOpenMPRuntimeGPU::emitWorkerLoop(CodeGenFunction &CGF,
1328 WorkerFunctionState &WST) {
1329 //
1330 // The workers enter this loop and wait for parallel work from the master.
1331 // When the master encounters a parallel region it sets up the work + variable
1332 // arguments, and wakes up the workers. The workers first check to see if
1333 // they are required for the parallel region, i.e., within the # of requested
1334 // parallel threads. The activated workers load the variable arguments and
1335 // execute the parallel work.
1336 //
1337
1338 CGBuilderTy &Bld = CGF.Builder;
1339
1340 llvm::BasicBlock *AwaitBB = CGF.createBasicBlock(".await.work");
1341 llvm::BasicBlock *SelectWorkersBB = CGF.createBasicBlock(".select.workers");
1342 llvm::BasicBlock *ExecuteBB = CGF.createBasicBlock(".execute.parallel");
1343 llvm::BasicBlock *TerminateBB = CGF.createBasicBlock(".terminate.parallel");
1344 llvm::BasicBlock *BarrierBB = CGF.createBasicBlock(".barrier.parallel");
1345 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".exit");
1346
1347 CGF.EmitBranch(AwaitBB);
1348
1349 // Workers wait for work from master.
1350 CGF.EmitBlock(AwaitBB);
1351 // Wait for parallel work
1352 syncCTAThreads(CGF);
1353
1354 Address WorkFn =
1355 CGF.CreateDefaultAlignTempAlloca(CGF.Int8PtrTy, /*Name=*/"work_fn");
1356 Address ExecStatus =
1357 CGF.CreateDefaultAlignTempAlloca(CGF.Int8Ty, /*Name=*/"exec_status");
1358 CGF.InitTempAlloca(ExecStatus, Bld.getInt8(/*C=*/0));
1359 CGF.InitTempAlloca(WorkFn, llvm::Constant::getNullValue(CGF.Int8PtrTy));
1360
1361 // TODO: Optimize runtime initialization and pass in correct value.
1362 llvm::Value *Args[] = {WorkFn.getPointer()};
1363 llvm::Value *Ret =
1364 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1365 CGM.getModule(), OMPRTL___kmpc_kernel_parallel),
1366 Args);
1367 Bld.CreateStore(Bld.CreateZExt(Ret, CGF.Int8Ty), ExecStatus);
1368
1369 // On termination condition (workid == 0), exit loop.
1370 llvm::Value *WorkID = Bld.CreateLoad(WorkFn);
1371 llvm::Value *ShouldTerminate = Bld.CreateIsNull(WorkID, "should_terminate");
1372 Bld.CreateCondBr(ShouldTerminate, ExitBB, SelectWorkersBB);
1373
1374 // Activate requested workers.
1375 CGF.EmitBlock(SelectWorkersBB);
1376 llvm::Value *IsActive =
1377 Bld.CreateIsNotNull(Bld.CreateLoad(ExecStatus), "is_active");
1378 Bld.CreateCondBr(IsActive, ExecuteBB, BarrierBB);
1379
1380 // Signal start of parallel region.
1381 CGF.EmitBlock(ExecuteBB);
1382 // Skip initialization.
1383 setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
1384
1385 // Process work items: outlined parallel functions.
1386 for (llvm::Function *W : Work) {
1387 // Try to match this outlined function.
1388 llvm::Value *ID = Bld.CreatePointerBitCastOrAddrSpaceCast(W, CGM.Int8PtrTy);
1389
1390 llvm::Value *WorkFnMatch =
1391 Bld.CreateICmpEQ(Bld.CreateLoad(WorkFn), ID, "work_match");
1392
1393 llvm::BasicBlock *ExecuteFNBB = CGF.createBasicBlock(".execute.fn");
1394 llvm::BasicBlock *CheckNextBB = CGF.createBasicBlock(".check.next");
1395 Bld.CreateCondBr(WorkFnMatch, ExecuteFNBB, CheckNextBB);
1396
1397 // Execute this outlined function.
1398 CGF.EmitBlock(ExecuteFNBB);
1399
1400 // Insert call to work function via shared wrapper. The shared
1401 // wrapper takes two arguments:
1402 // - the parallelism level;
1403 // - the thread ID;
1404 emitCall(CGF, WST.Loc, W,
1405 {Bld.getInt16(/*ParallelLevel=*/0), getThreadID(CGF, WST.Loc)});
1406
1407 // Go to end of parallel region.
1408 CGF.EmitBranch(TerminateBB);
1409
1410 CGF.EmitBlock(CheckNextBB);
1411 }
1412 // Default case: call to outlined function through pointer if the target
1413 // region makes a declare target call that may contain an orphaned parallel
1414 // directive.
1415 auto *ParallelFnTy =
1416 llvm::FunctionType::get(CGM.VoidTy, {CGM.Int16Ty, CGM.Int32Ty},
1417 /*isVarArg=*/false);
1418 llvm::Value *WorkFnCast =
1419 Bld.CreateBitCast(WorkID, ParallelFnTy->getPointerTo());
1420 // Insert call to work function via shared wrapper. The shared
1421 // wrapper takes two arguments:
1422 // - the parallelism level;
1423 // - the thread ID;
1424 emitCall(CGF, WST.Loc, {ParallelFnTy, WorkFnCast},
1425 {Bld.getInt16(/*ParallelLevel=*/0), getThreadID(CGF, WST.Loc)});
1426 // Go to end of parallel region.
1427 CGF.EmitBranch(TerminateBB);
1428
1429 // Signal end of parallel region.
1430 CGF.EmitBlock(TerminateBB);
1431 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1432 CGM.getModule(), OMPRTL___kmpc_kernel_end_parallel),
1433 llvm::None);
1434 CGF.EmitBranch(BarrierBB);
1435
1436 // All active and inactive workers wait at a barrier after parallel region.
1437 CGF.EmitBlock(BarrierBB);
1438 // Barrier after parallel region.
1439 syncCTAThreads(CGF);
1440 CGF.EmitBranch(AwaitBB);
1441
1442 // Exit target region.
1443 CGF.EmitBlock(ExitBB);
1444 // Skip initialization.
1445 clearLocThreadIdInsertPt(CGF);
1446}
1447
1448void CGOpenMPRuntimeGPU::createOffloadEntry(llvm::Constant *ID,
1449 llvm::Constant *Addr,
1450 uint64_t Size, int32_t,
1451 llvm::GlobalValue::LinkageTypes) {
1452 // TODO: Add support for global variables on the device after declare target
1453 // support.
1454 if (!isa<llvm::Function>(Addr))
1455 return;
1456 llvm::Module &M = CGM.getModule();
1457 llvm::LLVMContext &Ctx = CGM.getLLVMContext();
1458
1459 // Get "nvvm.annotations" metadata node
1460 llvm::NamedMDNode *MD = M.getOrInsertNamedMetadata("nvvm.annotations");
1461
1462 llvm::Metadata *MDVals[] = {
1463 llvm::ConstantAsMetadata::get(Addr), llvm::MDString::get(Ctx, "kernel"),
1464 llvm::ConstantAsMetadata::get(
1465 llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), 1))};
1466 // Append metadata to nvvm.annotations
1467 MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
1468}
1469
1470void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
1471 const OMPExecutableDirective &D, StringRef ParentName,
1472 llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
1473 bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
1474 if (!IsOffloadEntry) // Nothing to do.
1475 return;
1476
1477 assert(!ParentName.empty() && "Invalid target region parent name!")((!ParentName.empty() && "Invalid target region parent name!"
) ? static_cast<void> (0) : __assert_fail ("!ParentName.empty() && \"Invalid target region parent name!\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1477, __PRETTY_FUNCTION__))
;
1478
1479 bool Mode = supportsSPMDExecutionMode(CGM.getContext(), D);
1480 if (Mode)
1481 emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
1482 CodeGen);
1483 else
1484 emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
1485 CodeGen);
1486
1487 setPropertyExecutionMode(CGM, OutlinedFn->getName(), Mode);
1488}
1489
1490namespace {
1491LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE()using ::llvm::BitmaskEnumDetail::operator~; using ::llvm::BitmaskEnumDetail
::operator|; using ::llvm::BitmaskEnumDetail::operator&; using
::llvm::BitmaskEnumDetail::operator^; using ::llvm::BitmaskEnumDetail
::operator|=; using ::llvm::BitmaskEnumDetail::operator&=
; using ::llvm::BitmaskEnumDetail::operator^=
;
1492/// Enum for accesseing the reserved_2 field of the ident_t struct.
1493enum ModeFlagsTy : unsigned {
1494 /// Bit set to 1 when in SPMD mode.
1495 KMP_IDENT_SPMD_MODE = 0x01,
1496 /// Bit set to 1 when a simplified runtime is used.
1497 KMP_IDENT_SIMPLE_RT_MODE = 0x02,
1498 LLVM_MARK_AS_BITMASK_ENUM(/*LargestValue=*/KMP_IDENT_SIMPLE_RT_MODE)LLVM_BITMASK_LARGEST_ENUMERATOR = KMP_IDENT_SIMPLE_RT_MODE
1499};
1500
1501/// Special mode Undefined. Is the combination of Non-SPMD mode + SimpleRuntime.
1502static const ModeFlagsTy UndefinedMode =
1503 (~KMP_IDENT_SPMD_MODE) & KMP_IDENT_SIMPLE_RT_MODE;
1504} // anonymous namespace
1505
1506unsigned CGOpenMPRuntimeGPU::getDefaultLocationReserved2Flags() const {
1507 switch (getExecutionMode()) {
1508 case EM_SPMD:
1509 if (requiresFullRuntime())
1510 return KMP_IDENT_SPMD_MODE & (~KMP_IDENT_SIMPLE_RT_MODE);
1511 return KMP_IDENT_SPMD_MODE | KMP_IDENT_SIMPLE_RT_MODE;
1512 case EM_NonSPMD:
1513 assert(requiresFullRuntime() && "Expected full runtime.")((requiresFullRuntime() && "Expected full runtime.") ?
static_cast<void> (0) : __assert_fail ("requiresFullRuntime() && \"Expected full runtime.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1513, __PRETTY_FUNCTION__))
;
1514 return (~KMP_IDENT_SPMD_MODE) & (~KMP_IDENT_SIMPLE_RT_MODE);
1515 case EM_Unknown:
1516 return UndefinedMode;
1517 }
1518 llvm_unreachable("Unknown flags are requested.")::llvm::llvm_unreachable_internal("Unknown flags are requested."
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1518)
;
1519}
1520
1521CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
1522 : CGOpenMPRuntime(CGM, "_", "$") {
1523 if (!CGM.getLangOpts().OpenMPIsDevice)
1524 llvm_unreachable("OpenMP NVPTX can only handle device code.")::llvm::llvm_unreachable_internal("OpenMP NVPTX can only handle device code."
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1524)
;
1525}
1526
1527void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,
1528 ProcBindKind ProcBind,
1529 SourceLocation Loc) {
1530 // Do nothing in case of SPMD mode and L0 parallel.
1531 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
1532 return;
1533
1534 CGOpenMPRuntime::emitProcBindClause(CGF, ProcBind, Loc);
1535}
1536
1537void CGOpenMPRuntimeGPU::emitNumThreadsClause(CodeGenFunction &CGF,
1538 llvm::Value *NumThreads,
1539 SourceLocation Loc) {
1540 // Do nothing in case of SPMD mode and L0 parallel.
1541 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
1542 return;
1543
1544 CGOpenMPRuntime::emitNumThreadsClause(CGF, NumThreads, Loc);
1545}
1546
1547void CGOpenMPRuntimeGPU::emitNumTeamsClause(CodeGenFunction &CGF,
1548 const Expr *NumTeams,
1549 const Expr *ThreadLimit,
1550 SourceLocation Loc) {}
1551
1552llvm::Function *CGOpenMPRuntimeGPU::emitParallelOutlinedFunction(
1553 const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
1554 OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
1555 // Emit target region as a standalone region.
1556 class NVPTXPrePostActionTy : public PrePostActionTy {
1557 bool &IsInParallelRegion;
1558 bool PrevIsInParallelRegion;
1559
1560 public:
1561 NVPTXPrePostActionTy(bool &IsInParallelRegion)
1562 : IsInParallelRegion(IsInParallelRegion) {}
1563 void Enter(CodeGenFunction &CGF) override {
1564 PrevIsInParallelRegion = IsInParallelRegion;
1565 IsInParallelRegion = true;
1566 }
1567 void Exit(CodeGenFunction &CGF) override {
1568 IsInParallelRegion = PrevIsInParallelRegion;
1569 }
1570 } Action(IsInParallelRegion);
1571 CodeGen.setAction(Action);
1572 bool PrevIsInTTDRegion = IsInTTDRegion;
1573 IsInTTDRegion = false;
1574 bool PrevIsInTargetMasterThreadRegion = IsInTargetMasterThreadRegion;
1575 IsInTargetMasterThreadRegion = false;
1576 auto *OutlinedFun =
1577 cast<llvm::Function>(CGOpenMPRuntime::emitParallelOutlinedFunction(
1578 D, ThreadIDVar, InnermostKind, CodeGen));
1579 IsInTargetMasterThreadRegion = PrevIsInTargetMasterThreadRegion;
1580 IsInTTDRegion = PrevIsInTTDRegion;
1581 if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD &&
1582 !IsInParallelRegion) {
1583 llvm::Function *WrapperFun =
1584 createParallelDataSharingWrapper(OutlinedFun, D);
1585 WrapperFunctionsMap[OutlinedFun] = WrapperFun;
1586 }
1587
1588 return OutlinedFun;
1589}
1590
1591/// Get list of lastprivate variables from the teams distribute ... or
1592/// teams {distribute ...} directives.
1593static void
1594getDistributeLastprivateVars(ASTContext &Ctx, const OMPExecutableDirective &D,
1595 llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
1596 assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&((isOpenMPTeamsDirective(D.getDirectiveKind()) && "expected teams directive."
) ? static_cast<void> (0) : __assert_fail ("isOpenMPTeamsDirective(D.getDirectiveKind()) && \"expected teams directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1597, __PRETTY_FUNCTION__))
1597 "expected teams directive.")((isOpenMPTeamsDirective(D.getDirectiveKind()) && "expected teams directive."
) ? static_cast<void> (0) : __assert_fail ("isOpenMPTeamsDirective(D.getDirectiveKind()) && \"expected teams directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1597, __PRETTY_FUNCTION__))
;
1598 const OMPExecutableDirective *Dir = &D;
1599 if (!isOpenMPDistributeDirective(D.getDirectiveKind())) {
1600 if (const Stmt *S = CGOpenMPRuntime::getSingleCompoundChild(
1601 Ctx,
1602 D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(
1603 /*IgnoreCaptured=*/true))) {
1604 Dir = dyn_cast_or_null<OMPExecutableDirective>(S);
1605 if (Dir && !isOpenMPDistributeDirective(Dir->getDirectiveKind()))
1606 Dir = nullptr;
1607 }
1608 }
1609 if (!Dir)
1610 return;
1611 for (const auto *C : Dir->getClausesOfKind<OMPLastprivateClause>()) {
1612 for (const Expr *E : C->getVarRefs())
1613 Vars.push_back(getPrivateItem(E));
1614 }
1615}
1616
1617/// Get list of reduction variables from the teams ... directives.
1618static void
1619getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D,
1620 llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
1621 assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&((isOpenMPTeamsDirective(D.getDirectiveKind()) && "expected teams directive."
) ? static_cast<void> (0) : __assert_fail ("isOpenMPTeamsDirective(D.getDirectiveKind()) && \"expected teams directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1622, __PRETTY_FUNCTION__))
1622 "expected teams directive.")((isOpenMPTeamsDirective(D.getDirectiveKind()) && "expected teams directive."
) ? static_cast<void> (0) : __assert_fail ("isOpenMPTeamsDirective(D.getDirectiveKind()) && \"expected teams directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1622, __PRETTY_FUNCTION__))
;
1623 for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
1624 for (const Expr *E : C->privates())
1625 Vars.push_back(getPrivateItem(E));
1626 }
1627}
1628
1629llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(
1630 const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
1631 OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
1632 SourceLocation Loc = D.getBeginLoc();
1633
1634 const RecordDecl *GlobalizedRD = nullptr;
1635 llvm::SmallVector<const ValueDecl *, 4> LastPrivatesReductions;
1636 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
1637 unsigned WarpSize = CGM.getTarget().getGridValue(llvm::omp::GV_Warp_Size);
1638 // Globalize team reductions variable unconditionally in all modes.
1639 if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1640 getTeamsReductionVars(CGM.getContext(), D, LastPrivatesReductions);
1641 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
1642 getDistributeLastprivateVars(CGM.getContext(), D, LastPrivatesReductions);
1643 if (!LastPrivatesReductions.empty()) {
1644 GlobalizedRD = ::buildRecordForGlobalizedVars(
1645 CGM.getContext(), llvm::None, LastPrivatesReductions,
1646 MappedDeclsFields, WarpSize);
1647 }
1648 } else if (!LastPrivatesReductions.empty()) {
1649 assert(!TeamAndReductions.first &&((!TeamAndReductions.first && "Previous team declaration is not expected."
) ? static_cast<void> (0) : __assert_fail ("!TeamAndReductions.first && \"Previous team declaration is not expected.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1650, __PRETTY_FUNCTION__))
1650 "Previous team declaration is not expected.")((!TeamAndReductions.first && "Previous team declaration is not expected."
) ? static_cast<void> (0) : __assert_fail ("!TeamAndReductions.first && \"Previous team declaration is not expected.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1650, __PRETTY_FUNCTION__))
;
1651 TeamAndReductions.first = D.getCapturedStmt(OMPD_teams)->getCapturedDecl();
1652 std::swap(TeamAndReductions.second, LastPrivatesReductions);
1653 }
1654
1655 // Emit target region as a standalone region.
1656 class NVPTXPrePostActionTy : public PrePostActionTy {
1657 SourceLocation &Loc;
1658 const RecordDecl *GlobalizedRD;
1659 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1660 &MappedDeclsFields;
1661
1662 public:
1663 NVPTXPrePostActionTy(
1664 SourceLocation &Loc, const RecordDecl *GlobalizedRD,
1665 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1666 &MappedDeclsFields)
1667 : Loc(Loc), GlobalizedRD(GlobalizedRD),
1668 MappedDeclsFields(MappedDeclsFields) {}
1669 void Enter(CodeGenFunction &CGF) override {
1670 auto &Rt =
1671 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1672 if (GlobalizedRD) {
1673 auto I = Rt.FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;
1674 I->getSecond().GlobalRecord = GlobalizedRD;
1675 I->getSecond().MappedParams =
1676 std::make_unique<CodeGenFunction::OMPMapVars>();
1677 DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
1678 for (const auto &Pair : MappedDeclsFields) {
1679 assert(Pair.getFirst()->isCanonicalDecl() &&((Pair.getFirst()->isCanonicalDecl() && "Expected canonical declaration"
) ? static_cast<void> (0) : __assert_fail ("Pair.getFirst()->isCanonicalDecl() && \"Expected canonical declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1680, __PRETTY_FUNCTION__))
1680 "Expected canonical declaration")((Pair.getFirst()->isCanonicalDecl() && "Expected canonical declaration"
) ? static_cast<void> (0) : __assert_fail ("Pair.getFirst()->isCanonicalDecl() && \"Expected canonical declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1680, __PRETTY_FUNCTION__))
;
1681 Data.insert(std::make_pair(Pair.getFirst(),
1682 MappedVarData(Pair.getSecond(),
1683 /*IsOnePerTeam=*/true)));
1684 }
1685 }
1686 Rt.emitGenericVarsProlog(CGF, Loc);
1687 }
1688 void Exit(CodeGenFunction &CGF) override {
1689 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
1690 .emitGenericVarsEpilog(CGF);
1691 }
1692 } Action(Loc, GlobalizedRD, MappedDeclsFields);
1693 CodeGen.setAction(Action);
1694 llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(
1695 D, ThreadIDVar, InnermostKind, CodeGen);
1696
1697 return OutlinedFun;
1698}
1699
1700void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,
1701 SourceLocation Loc,
1702 bool WithSPMDCheck) {
1703 if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic &&
1704 getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1705 return;
1706
1707 CGBuilderTy &Bld = CGF.Builder;
1708
1709 const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
1710 if (I == FunctionGlobalizedDecls.end())
1711 return;
1712 if (const RecordDecl *GlobalizedVarsRecord = I->getSecond().GlobalRecord) {
1713 QualType GlobalRecTy = CGM.getContext().getRecordType(GlobalizedVarsRecord);
1714 QualType SecGlobalRecTy;
1715
1716 // Recover pointer to this function's global record. The runtime will
1717 // handle the specifics of the allocation of the memory.
1718 // Use actual memory size of the record including the padding
1719 // for alignment purposes.
1720 unsigned Alignment =
1721 CGM.getContext().getTypeAlignInChars(GlobalRecTy).getQuantity();
1722 unsigned GlobalRecordSize =
1723 CGM.getContext().getTypeSizeInChars(GlobalRecTy).getQuantity();
1724 GlobalRecordSize = llvm::alignTo(GlobalRecordSize, Alignment);
1725
1726 llvm::PointerType *GlobalRecPtrTy =
1727 CGF.ConvertTypeForMem(GlobalRecTy)->getPointerTo();
1728 llvm::Value *GlobalRecCastAddr;
1729 llvm::Value *IsTTD = nullptr;
1730 if (!IsInTTDRegion &&
1731 (WithSPMDCheck ||
1732 getExecutionMode() == CGOpenMPRuntimeGPU::EM_Unknown)) {
1733 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".exit");
1734 llvm::BasicBlock *SPMDBB = CGF.createBasicBlock(".spmd");
1735 llvm::BasicBlock *NonSPMDBB = CGF.createBasicBlock(".non-spmd");
1736 if (I->getSecond().SecondaryGlobalRecord.hasValue()) {
1737 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
1738 llvm::Value *ThreadID = getThreadID(CGF, Loc);
1739 llvm::Value *PL = CGF.EmitRuntimeCall(
1740 OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
1741 OMPRTL___kmpc_parallel_level),
1742 {RTLoc, ThreadID});
1743 IsTTD = Bld.CreateIsNull(PL);
1744 }
1745 llvm::Value *IsSPMD = Bld.CreateIsNotNull(
1746 CGF.EmitNounwindRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1747 CGM.getModule(), OMPRTL___kmpc_is_spmd_exec_mode)));
1748 Bld.CreateCondBr(IsSPMD, SPMDBB, NonSPMDBB);
1749 // There is no need to emit line number for unconditional branch.
1750 (void)ApplyDebugLocation::CreateEmpty(CGF);
1751 CGF.EmitBlock(SPMDBB);
1752 Address RecPtr = Address(llvm::ConstantPointerNull::get(GlobalRecPtrTy),
1753 CharUnits::fromQuantity(Alignment));
1754 CGF.EmitBranch(ExitBB);
1755 // There is no need to emit line number for unconditional branch.
1756 (void)ApplyDebugLocation::CreateEmpty(CGF);
1757 CGF.EmitBlock(NonSPMDBB);
1758 llvm::Value *Size = llvm::ConstantInt::get(CGM.SizeTy, GlobalRecordSize);
1759 if (const RecordDecl *SecGlobalizedVarsRecord =
1760 I->getSecond().SecondaryGlobalRecord.getValueOr(nullptr)) {
1761 SecGlobalRecTy =
1762 CGM.getContext().getRecordType(SecGlobalizedVarsRecord);
1763
1764 // Recover pointer to this function's global record. The runtime will
1765 // handle the specifics of the allocation of the memory.
1766 // Use actual memory size of the record including the padding
1767 // for alignment purposes.
1768 unsigned Alignment =
1769 CGM.getContext().getTypeAlignInChars(SecGlobalRecTy).getQuantity();
1770 unsigned GlobalRecordSize =
1771 CGM.getContext().getTypeSizeInChars(SecGlobalRecTy).getQuantity();
1772 GlobalRecordSize = llvm::alignTo(GlobalRecordSize, Alignment);
1773 Size = Bld.CreateSelect(
1774 IsTTD, llvm::ConstantInt::get(CGM.SizeTy, GlobalRecordSize), Size);
1775 }
1776 // TODO: allow the usage of shared memory to be controlled by
1777 // the user, for now, default to global.
1778 llvm::Value *GlobalRecordSizeArg[] = {
1779 Size, CGF.Builder.getInt16(/*UseSharedMemory=*/0)};
1780 llvm::Value *GlobalRecValue = CGF.EmitRuntimeCall(
1781 OMPBuilder.getOrCreateRuntimeFunction(
1782 CGM.getModule(), OMPRTL___kmpc_data_sharing_coalesced_push_stack),
1783 GlobalRecordSizeArg);
1784 GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1785 GlobalRecValue, GlobalRecPtrTy);
1786 CGF.EmitBlock(ExitBB);
1787 auto *Phi = Bld.CreatePHI(GlobalRecPtrTy,
1788 /*NumReservedValues=*/2, "_select_stack");
1789 Phi->addIncoming(RecPtr.getPointer(), SPMDBB);
1790 Phi->addIncoming(GlobalRecCastAddr, NonSPMDBB);
1791 GlobalRecCastAddr = Phi;
1792 I->getSecond().GlobalRecordAddr = Phi;
1793 I->getSecond().IsInSPMDModeFlag = IsSPMD;
1794 } else if (!CGM.getLangOpts().OpenMPCUDATargetParallel && IsInTTDRegion) {
1795 assert(GlobalizedRecords.back().Records.size() < 2 &&((GlobalizedRecords.back().Records.size() < 2 && "Expected less than 2 globalized records: one for target and one "
"for teams.") ? static_cast<void> (0) : __assert_fail (
"GlobalizedRecords.back().Records.size() < 2 && \"Expected less than 2 globalized records: one for target and one \" \"for teams.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1797, __PRETTY_FUNCTION__))
1796 "Expected less than 2 globalized records: one for target and one "((GlobalizedRecords.back().Records.size() < 2 && "Expected less than 2 globalized records: one for target and one "
"for teams.") ? static_cast<void> (0) : __assert_fail (
"GlobalizedRecords.back().Records.size() < 2 && \"Expected less than 2 globalized records: one for target and one \" \"for teams.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1797, __PRETTY_FUNCTION__))
1797 "for teams.")((GlobalizedRecords.back().Records.size() < 2 && "Expected less than 2 globalized records: one for target and one "
"for teams.") ? static_cast<void> (0) : __assert_fail (
"GlobalizedRecords.back().Records.size() < 2 && \"Expected less than 2 globalized records: one for target and one \" \"for teams.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1797, __PRETTY_FUNCTION__))
;
1798 unsigned Offset = 0;
1799 for (const RecordDecl *RD : GlobalizedRecords.back().Records) {
1800 QualType RDTy = CGM.getContext().getRecordType(RD);
1801 unsigned Alignment =
1802 CGM.getContext().getTypeAlignInChars(RDTy).getQuantity();
1803 unsigned Size = CGM.getContext().getTypeSizeInChars(RDTy).getQuantity();
1804 Offset =
1805 llvm::alignTo(llvm::alignTo(Offset, Alignment) + Size, Alignment);
1806 }
1807 unsigned Alignment =
1808 CGM.getContext().getTypeAlignInChars(GlobalRecTy).getQuantity();
1809 Offset = llvm::alignTo(Offset, Alignment);
1810 GlobalizedRecords.back().Records.push_back(GlobalizedVarsRecord);
1811 ++GlobalizedRecords.back().RegionCounter;
1812 if (GlobalizedRecords.back().Records.size() == 1) {
1813 assert(KernelStaticGlobalized &&((KernelStaticGlobalized && "Kernel static pointer must be initialized already."
) ? static_cast<void> (0) : __assert_fail ("KernelStaticGlobalized && \"Kernel static pointer must be initialized already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1814, __PRETTY_FUNCTION__))
1814 "Kernel static pointer must be initialized already.")((KernelStaticGlobalized && "Kernel static pointer must be initialized already."
) ? static_cast<void> (0) : __assert_fail ("KernelStaticGlobalized && \"Kernel static pointer must be initialized already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1814, __PRETTY_FUNCTION__))
;
1815 auto *UseSharedMemory = new llvm::GlobalVariable(
1816 CGM.getModule(), CGM.Int16Ty, /*isConstant=*/true,
1817 llvm::GlobalValue::InternalLinkage, nullptr,
1818 "_openmp_static_kernel$is_shared");
1819 UseSharedMemory->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
1820 QualType Int16Ty = CGM.getContext().getIntTypeForBitwidth(
1821 /*DestWidth=*/16, /*Signed=*/0);
1822 llvm::Value *IsInSharedMemory = CGF.EmitLoadOfScalar(
1823 Address(UseSharedMemory,
1824 CGM.getContext().getTypeAlignInChars(Int16Ty)),
1825 /*Volatile=*/false, Int16Ty, Loc);
1826 auto *StaticGlobalized = new llvm::GlobalVariable(
1827 CGM.getModule(), CGM.Int8Ty, /*isConstant=*/false,
1828 llvm::GlobalValue::CommonLinkage, nullptr);
1829 auto *RecSize = new llvm::GlobalVariable(
1830 CGM.getModule(), CGM.SizeTy, /*isConstant=*/true,
1831 llvm::GlobalValue::InternalLinkage, nullptr,
1832 "_openmp_static_kernel$size");
1833 RecSize->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
1834 llvm::Value *Ld = CGF.EmitLoadOfScalar(
1835 Address(RecSize, CGM.getSizeAlign()), /*Volatile=*/false,
1836 CGM.getContext().getSizeType(), Loc);
1837 llvm::Value *ResAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1838 KernelStaticGlobalized, CGM.VoidPtrPtrTy);
1839 llvm::Value *GlobalRecordSizeArg[] = {
1840 llvm::ConstantInt::get(
1841 CGM.Int16Ty,
1842 getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD ? 1 : 0),
1843 StaticGlobalized, Ld, IsInSharedMemory, ResAddr};
1844 CGF.EmitRuntimeCall(
1845 OMPBuilder.getOrCreateRuntimeFunction(
1846 CGM.getModule(), OMPRTL___kmpc_get_team_static_memory),
1847 GlobalRecordSizeArg);
1848 GlobalizedRecords.back().Buffer = StaticGlobalized;
1849 GlobalizedRecords.back().RecSize = RecSize;
1850 GlobalizedRecords.back().UseSharedMemory = UseSharedMemory;
1851 GlobalizedRecords.back().Loc = Loc;
1852 }
1853 assert(KernelStaticGlobalized && "Global address must be set already.")((KernelStaticGlobalized && "Global address must be set already."
) ? static_cast<void> (0) : __assert_fail ("KernelStaticGlobalized && \"Global address must be set already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1853, __PRETTY_FUNCTION__))
;
1854 Address FrameAddr = CGF.EmitLoadOfPointer(
1855 Address(KernelStaticGlobalized, CGM.getPointerAlign()),
1856 CGM.getContext()
1857 .getPointerType(CGM.getContext().VoidPtrTy)
1858 .castAs<PointerType>());
1859 llvm::Value *GlobalRecValue =
1860 Bld.CreateConstInBoundsGEP(FrameAddr, Offset).getPointer();
1861 I->getSecond().GlobalRecordAddr = GlobalRecValue;
1862 I->getSecond().IsInSPMDModeFlag = nullptr;
1863 GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1864 GlobalRecValue, CGF.ConvertTypeForMem(GlobalRecTy)->getPointerTo());
1865 } else {
1866 // TODO: allow the usage of shared memory to be controlled by
1867 // the user, for now, default to global.
1868 bool UseSharedMemory =
1869 IsInTTDRegion && GlobalRecordSize <= SharedMemorySize;
1870 llvm::Value *GlobalRecordSizeArg[] = {
1871 llvm::ConstantInt::get(CGM.SizeTy, GlobalRecordSize),
1872 CGF.Builder.getInt16(UseSharedMemory ? 1 : 0)};
1873 llvm::Value *GlobalRecValue = CGF.EmitRuntimeCall(
1874 OMPBuilder.getOrCreateRuntimeFunction(
1875 CGM.getModule(),
1876 IsInTTDRegion ? OMPRTL___kmpc_data_sharing_push_stack
1877 : OMPRTL___kmpc_data_sharing_coalesced_push_stack),
1878 GlobalRecordSizeArg);
1879 GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1880 GlobalRecValue, GlobalRecPtrTy);
1881 I->getSecond().GlobalRecordAddr = GlobalRecValue;
1882 I->getSecond().IsInSPMDModeFlag = nullptr;
1883 }
1884 LValue Base =
1885 CGF.MakeNaturalAlignPointeeAddrLValue(GlobalRecCastAddr, GlobalRecTy);
1886
1887 // Emit the "global alloca" which is a GEP from the global declaration
1888 // record using the pointer returned by the runtime.
1889 LValue SecBase;
1890 decltype(I->getSecond().LocalVarData)::const_iterator SecIt;
1891 if (IsTTD) {
1892 SecIt = I->getSecond().SecondaryLocalVarData->begin();
1893 llvm::PointerType *SecGlobalRecPtrTy =
1894 CGF.ConvertTypeForMem(SecGlobalRecTy)->getPointerTo();
1895 SecBase = CGF.MakeNaturalAlignPointeeAddrLValue(
1896 Bld.CreatePointerBitCastOrAddrSpaceCast(
1897 I->getSecond().GlobalRecordAddr, SecGlobalRecPtrTy),
1898 SecGlobalRecTy);
1899 }
1900 for (auto &Rec : I->getSecond().LocalVarData) {
1901 bool EscapedParam = I->getSecond().EscapedParameters.count(Rec.first);
1902 llvm::Value *ParValue;
1903 if (EscapedParam) {
1904 const auto *VD = cast<VarDecl>(Rec.first);
1905 LValue ParLVal =
1906 CGF.MakeAddrLValue(CGF.GetAddrOfLocalVar(VD), VD->getType());
1907 ParValue = CGF.EmitLoadOfScalar(ParLVal, Loc);
1908 }
1909 LValue VarAddr = CGF.EmitLValueForField(Base, Rec.second.FD);
1910 // Emit VarAddr basing on lane-id if required.
1911 QualType VarTy;
1912 if (Rec.second.IsOnePerTeam) {
1913 VarTy = Rec.second.FD->getType();
1914 } else {
1915 Address Addr = VarAddr.getAddress(CGF);
1916 llvm::Value *Ptr = CGF.Builder.CreateInBoundsGEP(
1917 Addr.getElementType(), Addr.getPointer(),
1918 {Bld.getInt32(0), getNVPTXLaneID(CGF)});
1919 VarTy =
1920 Rec.second.FD->getType()->castAsArrayTypeUnsafe()->getElementType();
1921 VarAddr = CGF.MakeAddrLValue(
1922 Address(Ptr, CGM.getContext().getDeclAlign(Rec.first)), VarTy,
1923 AlignmentSource::Decl);
1924 }
1925 Rec.second.PrivateAddr = VarAddr.getAddress(CGF);
1926 if (!IsInTTDRegion &&
1927 (WithSPMDCheck ||
1928 getExecutionMode() == CGOpenMPRuntimeGPU::EM_Unknown)) {
1929 assert(I->getSecond().IsInSPMDModeFlag &&((I->getSecond().IsInSPMDModeFlag && "Expected unknown execution mode or required SPMD check."
) ? static_cast<void> (0) : __assert_fail ("I->getSecond().IsInSPMDModeFlag && \"Expected unknown execution mode or required SPMD check.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1930, __PRETTY_FUNCTION__))
1930 "Expected unknown execution mode or required SPMD check.")((I->getSecond().IsInSPMDModeFlag && "Expected unknown execution mode or required SPMD check."
) ? static_cast<void> (0) : __assert_fail ("I->getSecond().IsInSPMDModeFlag && \"Expected unknown execution mode or required SPMD check.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1930, __PRETTY_FUNCTION__))
;
1931 if (IsTTD) {
1932 assert(SecIt->second.IsOnePerTeam &&((SecIt->second.IsOnePerTeam && "Secondary glob data must be one per team."
) ? static_cast<void> (0) : __assert_fail ("SecIt->second.IsOnePerTeam && \"Secondary glob data must be one per team.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1933, __PRETTY_FUNCTION__))
1933 "Secondary glob data must be one per team.")((SecIt->second.IsOnePerTeam && "Secondary glob data must be one per team."
) ? static_cast<void> (0) : __assert_fail ("SecIt->second.IsOnePerTeam && \"Secondary glob data must be one per team.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 1933, __PRETTY_FUNCTION__))
;
1934 LValue SecVarAddr = CGF.EmitLValueForField(SecBase, SecIt->second.FD);
1935 VarAddr.setAddress(
1936 Address(Bld.CreateSelect(IsTTD, SecVarAddr.getPointer(CGF),
1937 VarAddr.getPointer(CGF)),
1938 VarAddr.getAlignment()));
1939 Rec.second.PrivateAddr = VarAddr.getAddress(CGF);
1940 }
1941 Address GlobalPtr = Rec.second.PrivateAddr;
1942 Address LocalAddr = CGF.CreateMemTemp(VarTy, Rec.second.FD->getName());
1943 Rec.second.PrivateAddr = Address(
1944 Bld.CreateSelect(I->getSecond().IsInSPMDModeFlag,
1945 LocalAddr.getPointer(), GlobalPtr.getPointer()),
1946 LocalAddr.getAlignment());
1947 }
1948 if (EscapedParam) {
1949 const auto *VD = cast<VarDecl>(Rec.first);
1950 CGF.EmitStoreOfScalar(ParValue, VarAddr);
1951 I->getSecond().MappedParams->setVarAddr(CGF, VD,
1952 VarAddr.getAddress(CGF));
1953 }
1954 if (IsTTD)
1955 ++SecIt;
1956 }
1957 }
1958 for (const ValueDecl *VD : I->getSecond().EscapedVariableLengthDecls) {
1959 // Recover pointer to this function's global record. The runtime will
1960 // handle the specifics of the allocation of the memory.
1961 // Use actual memory size of the record including the padding
1962 // for alignment purposes.
1963 CGBuilderTy &Bld = CGF.Builder;
1964 llvm::Value *Size = CGF.getTypeSize(VD->getType());
1965 CharUnits Align = CGM.getContext().getDeclAlign(VD);
1966 Size = Bld.CreateNUWAdd(
1967 Size, llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity() - 1));
1968 llvm::Value *AlignVal =
1969 llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity());
1970 Size = Bld.CreateUDiv(Size, AlignVal);
1971 Size = Bld.CreateNUWMul(Size, AlignVal);
1972 // TODO: allow the usage of shared memory to be controlled by
1973 // the user, for now, default to global.
1974 llvm::Value *GlobalRecordSizeArg[] = {
1975 Size, CGF.Builder.getInt16(/*UseSharedMemory=*/0)};
1976 llvm::Value *GlobalRecValue = CGF.EmitRuntimeCall(
1977 OMPBuilder.getOrCreateRuntimeFunction(
1978 CGM.getModule(), OMPRTL___kmpc_data_sharing_coalesced_push_stack),
1979 GlobalRecordSizeArg);
1980 llvm::Value *GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1981 GlobalRecValue, CGF.ConvertTypeForMem(VD->getType())->getPointerTo());
1982 LValue Base = CGF.MakeAddrLValue(GlobalRecCastAddr, VD->getType(),
1983 CGM.getContext().getDeclAlign(VD),
1984 AlignmentSource::Decl);
1985 I->getSecond().MappedParams->setVarAddr(CGF, cast<VarDecl>(VD),
1986 Base.getAddress(CGF));
1987 I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(GlobalRecValue);
1988 }
1989 I->getSecond().MappedParams->apply(CGF);
1990}
1991
1992void CGOpenMPRuntimeGPU::emitGenericVarsEpilog(CodeGenFunction &CGF,
1993 bool WithSPMDCheck) {
1994 if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic &&
1995 getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1996 return;
1997
1998 const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
1999 if (I != FunctionGlobalizedDecls.end()) {
2000 I->getSecond().MappedParams->restore(CGF);
2001 if (!CGF.HaveInsertPoint())
2002 return;
2003 for (llvm::Value *Addr :
2004 llvm::reverse(I->getSecond().EscapedVariableLengthDeclsAddrs)) {
2005 CGF.EmitRuntimeCall(
2006 OMPBuilder.getOrCreateRuntimeFunction(
2007 CGM.getModule(), OMPRTL___kmpc_data_sharing_pop_stack),
2008 Addr);
2009 }
2010 if (I->getSecond().GlobalRecordAddr) {
2011 if (!IsInTTDRegion &&
2012 (WithSPMDCheck ||
2013 getExecutionMode() == CGOpenMPRuntimeGPU::EM_Unknown)) {
2014 CGBuilderTy &Bld = CGF.Builder;
2015 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".exit");
2016 llvm::BasicBlock *NonSPMDBB = CGF.createBasicBlock(".non-spmd");
2017 Bld.CreateCondBr(I->getSecond().IsInSPMDModeFlag, ExitBB, NonSPMDBB);
2018 // There is no need to emit line number for unconditional branch.
2019 (void)ApplyDebugLocation::CreateEmpty(CGF);
2020 CGF.EmitBlock(NonSPMDBB);
2021 CGF.EmitRuntimeCall(
2022 OMPBuilder.getOrCreateRuntimeFunction(
2023 CGM.getModule(), OMPRTL___kmpc_data_sharing_pop_stack),
2024 CGF.EmitCastToVoidPtr(I->getSecond().GlobalRecordAddr));
2025 CGF.EmitBlock(ExitBB);
2026 } else if (!CGM.getLangOpts().OpenMPCUDATargetParallel && IsInTTDRegion) {
2027 assert(GlobalizedRecords.back().RegionCounter > 0 &&((GlobalizedRecords.back().RegionCounter > 0 && "region counter must be > 0."
) ? static_cast<void> (0) : __assert_fail ("GlobalizedRecords.back().RegionCounter > 0 && \"region counter must be > 0.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2028, __PRETTY_FUNCTION__))
2028 "region counter must be > 0.")((GlobalizedRecords.back().RegionCounter > 0 && "region counter must be > 0."
) ? static_cast<void> (0) : __assert_fail ("GlobalizedRecords.back().RegionCounter > 0 && \"region counter must be > 0.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2028, __PRETTY_FUNCTION__))
;
2029 --GlobalizedRecords.back().RegionCounter;
2030 // Emit the restore function only in the target region.
2031 if (GlobalizedRecords.back().RegionCounter == 0) {
2032 QualType Int16Ty = CGM.getContext().getIntTypeForBitwidth(
2033 /*DestWidth=*/16, /*Signed=*/0);
2034 llvm::Value *IsInSharedMemory = CGF.EmitLoadOfScalar(
2035 Address(GlobalizedRecords.back().UseSharedMemory,
2036 CGM.getContext().getTypeAlignInChars(Int16Ty)),
2037 /*Volatile=*/false, Int16Ty, GlobalizedRecords.back().Loc);
2038 llvm::Value *Args[] = {
2039 llvm::ConstantInt::get(
2040 CGM.Int16Ty,
2041 getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD ? 1 : 0),
2042 IsInSharedMemory};
2043 CGF.EmitRuntimeCall(
2044 OMPBuilder.getOrCreateRuntimeFunction(
2045 CGM.getModule(), OMPRTL___kmpc_restore_team_static_memory),
2046 Args);
2047 }
2048 } else {
2049 CGF.EmitRuntimeCall(
2050 OMPBuilder.getOrCreateRuntimeFunction(
2051 CGM.getModule(), OMPRTL___kmpc_data_sharing_pop_stack),
2052 I->getSecond().GlobalRecordAddr);
2053 }
2054 }
2055 }
2056}
2057
2058void CGOpenMPRuntimeGPU::emitTeamsCall(CodeGenFunction &CGF,
2059 const OMPExecutableDirective &D,
2060 SourceLocation Loc,
2061 llvm::Function *OutlinedFn,
2062 ArrayRef<llvm::Value *> CapturedVars) {
2063 if (!CGF.HaveInsertPoint())
2064 return;
2065
2066 Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
2067 /*Name=*/".zero.addr");
2068 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
2069 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
2070 OutlinedFnArgs.push_back(emitThreadIDAddress(CGF, Loc).getPointer());
2071 OutlinedFnArgs.push_back(ZeroAddr.getPointer());
2072 OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
2073 emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, OutlinedFnArgs);
2074}
2075
2076void CGOpenMPRuntimeGPU::emitParallelCall(
2077 CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *OutlinedFn,
2078 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
2079 if (!CGF.HaveInsertPoint())
2080 return;
2081
2082 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
2083 emitSPMDParallelCall(CGF, Loc, OutlinedFn, CapturedVars, IfCond);
2084 else
2085 emitNonSPMDParallelCall(CGF, Loc, OutlinedFn, CapturedVars, IfCond);
2086}
2087
2088void CGOpenMPRuntimeGPU::emitNonSPMDParallelCall(
2089 CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
2090 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
2091 llvm::Function *Fn = cast<llvm::Function>(OutlinedFn);
2092
2093 // Force inline this outlined function at its call site.
2094 Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
2095
2096 Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
2097 /*Name=*/".zero.addr");
2098 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
2099 // ThreadId for serialized parallels is 0.
2100 Address ThreadIDAddr = ZeroAddr;
2101 auto &&CodeGen = [this, Fn, CapturedVars, Loc, &ThreadIDAddr](
2102 CodeGenFunction &CGF, PrePostActionTy &Action) {
2103 Action.Enter(CGF);
2104
2105 Address ZeroAddr =
2106 CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
2107 /*Name=*/".bound.zero.addr");
2108 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
2109 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
2110 OutlinedFnArgs.push_back(ThreadIDAddr.getPointer());
2111 OutlinedFnArgs.push_back(ZeroAddr.getPointer());
2112 OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
2113 emitOutlinedFunctionCall(CGF, Loc, Fn, OutlinedFnArgs);
2114 };
2115 auto &&SeqGen = [this, &CodeGen, Loc](CodeGenFunction &CGF,
2116 PrePostActionTy &) {
2117
2118 RegionCodeGenTy RCG(CodeGen);
2119 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
2120 llvm::Value *ThreadID = getThreadID(CGF, Loc);
2121 llvm::Value *Args[] = {RTLoc, ThreadID};
2122
2123 NVPTXActionTy Action(
2124 OMPBuilder.getOrCreateRuntimeFunction(
2125 CGM.getModule(), OMPRTL___kmpc_serialized_parallel),
2126 Args,
2127 OMPBuilder.getOrCreateRuntimeFunction(
2128 CGM.getModule(), OMPRTL___kmpc_end_serialized_parallel),
2129 Args);
2130 RCG.setAction(Action);
2131 RCG(CGF);
2132 };
2133
2134 auto &&L0ParallelGen = [this, CapturedVars, Fn](CodeGenFunction &CGF,
2135 PrePostActionTy &Action) {
2136 CGBuilderTy &Bld = CGF.Builder;
2137 llvm::Function *WFn = WrapperFunctionsMap[Fn];
2138 assert(WFn && "Wrapper function does not exist!")((WFn && "Wrapper function does not exist!") ? static_cast
<void> (0) : __assert_fail ("WFn && \"Wrapper function does not exist!\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2138, __PRETTY_FUNCTION__))
;
2139 llvm::Value *ID = Bld.CreateBitOrPointerCast(WFn, CGM.Int8PtrTy);
2140
2141 // Prepare for parallel region. Indicate the outlined function.
2142 llvm::Value *Args[] = {ID};
2143 CGF.EmitRuntimeCall(
2144 OMPBuilder.getOrCreateRuntimeFunction(
2145 CGM.getModule(), OMPRTL___kmpc_kernel_prepare_parallel),
2146 Args);
2147
2148 // Create a private scope that will globalize the arguments
2149 // passed from the outside of the target region.
2150 CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF);
2151
2152 // There's something to share.
2153 if (!CapturedVars.empty()) {
2154 // Prepare for parallel region. Indicate the outlined function.
2155 Address SharedArgs =
2156 CGF.CreateDefaultAlignTempAlloca(CGF.VoidPtrPtrTy, "shared_arg_refs");
2157 llvm::Value *SharedArgsPtr = SharedArgs.getPointer();
2158
2159 llvm::Value *DataSharingArgs[] = {
2160 SharedArgsPtr,
2161 llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
2162 CGF.EmitRuntimeCall(
2163 OMPBuilder.getOrCreateRuntimeFunction(
2164 CGM.getModule(), OMPRTL___kmpc_begin_sharing_variables),
2165 DataSharingArgs);
2166
2167 // Store variable address in a list of references to pass to workers.
2168 unsigned Idx = 0;
2169 ASTContext &Ctx = CGF.getContext();
2170 Address SharedArgListAddress = CGF.EmitLoadOfPointer(
2171 SharedArgs, Ctx.getPointerType(Ctx.getPointerType(Ctx.VoidPtrTy))
2172 .castAs<PointerType>());
2173 for (llvm::Value *V : CapturedVars) {
2174 Address Dst = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);
2175 llvm::Value *PtrV;
2176 if (V->getType()->isIntegerTy())
2177 PtrV = Bld.CreateIntToPtr(V, CGF.VoidPtrTy);
2178 else
2179 PtrV = Bld.CreatePointerBitCastOrAddrSpaceCast(V, CGF.VoidPtrTy);
2180 CGF.EmitStoreOfScalar(PtrV, Dst, /*Volatile=*/false,
2181 Ctx.getPointerType(Ctx.VoidPtrTy));
2182 ++Idx;
2183 }
2184 }
2185
2186 // Activate workers. This barrier is used by the master to signal
2187 // work for the workers.
2188 syncCTAThreads(CGF);
2189
2190 // OpenMP [2.5, Parallel Construct, p.49]
2191 // There is an implied barrier at the end of a parallel region. After the
2192 // end of a parallel region, only the master thread of the team resumes
2193 // execution of the enclosing task region.
2194 //
2195 // The master waits at this barrier until all workers are done.
2196 syncCTAThreads(CGF);
2197
2198 if (!CapturedVars.empty())
2199 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
2200 CGM.getModule(), OMPRTL___kmpc_end_sharing_variables));
2201
2202 // Remember for post-processing in worker loop.
2203 Work.emplace_back(WFn);
2204 };
2205
2206 auto &&LNParallelGen = [this, Loc, &SeqGen, &L0ParallelGen](
2207 CodeGenFunction &CGF, PrePostActionTy &Action) {
2208 if (IsInParallelRegion) {
2209 SeqGen(CGF, Action);
2210 } else if (IsInTargetMasterThreadRegion) {
2211 L0ParallelGen(CGF, Action);
2212 } else {
2213 // Check for master and then parallelism:
2214 // if (__kmpc_is_spmd_exec_mode() || __kmpc_parallel_level(loc, gtid)) {
2215 // Serialized execution.
2216 // } else {
2217 // Worker call.
2218 // }
2219 CGBuilderTy &Bld = CGF.Builder;
2220 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".exit");
2221 llvm::BasicBlock *SeqBB = CGF.createBasicBlock(".sequential");
2222 llvm::BasicBlock *ParallelCheckBB = CGF.createBasicBlock(".parcheck");
2223 llvm::BasicBlock *MasterBB = CGF.createBasicBlock(".master");
2224 llvm::Value *IsSPMD = Bld.CreateIsNotNull(
2225 CGF.EmitNounwindRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
2226 CGM.getModule(), OMPRTL___kmpc_is_spmd_exec_mode)));
2227 Bld.CreateCondBr(IsSPMD, SeqBB, ParallelCheckBB);
2228 // There is no need to emit line number for unconditional branch.
2229 (void)ApplyDebugLocation::CreateEmpty(CGF);
2230 CGF.EmitBlock(ParallelCheckBB);
2231 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
2232 llvm::Value *ThreadID = getThreadID(CGF, Loc);
2233 llvm::Value *PL = CGF.EmitRuntimeCall(
2234 OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
2235 OMPRTL___kmpc_parallel_level),
2236 {RTLoc, ThreadID});
2237 llvm::Value *Res = Bld.CreateIsNotNull(PL);
2238 Bld.CreateCondBr(Res, SeqBB, MasterBB);
2239 CGF.EmitBlock(SeqBB);
2240 SeqGen(CGF, Action);
2241 CGF.EmitBranch(ExitBB);
2242 // There is no need to emit line number for unconditional branch.
2243 (void)ApplyDebugLocation::CreateEmpty(CGF);
2244 CGF.EmitBlock(MasterBB);
2245 L0ParallelGen(CGF, Action);
2246 CGF.EmitBranch(ExitBB);
2247 // There is no need to emit line number for unconditional branch.
2248 (void)ApplyDebugLocation::CreateEmpty(CGF);
2249 // Emit the continuation block for code after the if.
2250 CGF.EmitBlock(ExitBB, /*IsFinished=*/true);
2251 }
2252 };
2253
2254 if (IfCond) {
2255 emitIfClause(CGF, IfCond, LNParallelGen, SeqGen);
2256 } else {
2257 CodeGenFunction::RunCleanupsScope Scope(CGF);
2258 RegionCodeGenTy ThenRCG(LNParallelGen);
2259 ThenRCG(CGF);
2260 }
2261}
2262
2263void CGOpenMPRuntimeGPU::emitSPMDParallelCall(
2264 CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *OutlinedFn,
2265 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
2266 // Just call the outlined function to execute the parallel region.
2267 // OutlinedFn(&GTid, &zero, CapturedStruct);
2268 //
2269 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
2270
2271 Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
2272 /*Name=*/".zero.addr");
2273 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
2274 // ThreadId for serialized parallels is 0.
2275 Address ThreadIDAddr = ZeroAddr;
2276 auto &&CodeGen = [this, OutlinedFn, CapturedVars, Loc, &ThreadIDAddr](
2277 CodeGenFunction &CGF, PrePostActionTy &Action) {
2278 Action.Enter(CGF);
2279
2280 Address ZeroAddr =
2281 CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
2282 /*Name=*/".bound.zero.addr");
2283 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
2284 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
2285 OutlinedFnArgs.push_back(ThreadIDAddr.getPointer());
2286 OutlinedFnArgs.push_back(ZeroAddr.getPointer());
2287 OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
2288 emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, OutlinedFnArgs);
2289 };
2290 auto &&SeqGen = [this, &CodeGen, Loc](CodeGenFunction &CGF,
2291 PrePostActionTy &) {
2292
2293 RegionCodeGenTy RCG(CodeGen);
2294 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
2295 llvm::Value *ThreadID = getThreadID(CGF, Loc);
2296 llvm::Value *Args[] = {RTLoc, ThreadID};
2297
2298 NVPTXActionTy Action(
2299 OMPBuilder.getOrCreateRuntimeFunction(
2300 CGM.getModule(), OMPRTL___kmpc_serialized_parallel),
2301 Args,
2302 OMPBuilder.getOrCreateRuntimeFunction(
2303 CGM.getModule(), OMPRTL___kmpc_end_serialized_parallel),
2304 Args);
2305 RCG.setAction(Action);
2306 RCG(CGF);
2307 };
2308
2309 if (IsInTargetMasterThreadRegion) {
2310 // In the worker need to use the real thread id.
2311 ThreadIDAddr = emitThreadIDAddress(CGF, Loc);
2312 RegionCodeGenTy RCG(CodeGen);
2313 RCG(CGF);
2314 } else {
2315 // If we are not in the target region, it is definitely L2 parallelism or
2316 // more, because for SPMD mode we always has L1 parallel level, sowe don't
2317 // need to check for orphaned directives.
2318 RegionCodeGenTy RCG(SeqGen);
2319 RCG(CGF);
2320 }
2321}
2322
2323void CGOpenMPRuntimeGPU::syncCTAThreads(CodeGenFunction &CGF) {
2324 // Always emit simple barriers!
2325 if (!CGF.HaveInsertPoint())
2326 return;
2327 // Build call __kmpc_barrier_simple_spmd(nullptr, 0);
2328 // This function does not use parameters, so we can emit just default values.
2329 llvm::Value *Args[] = {
2330 llvm::ConstantPointerNull::get(
2331 cast<llvm::PointerType>(getIdentTyPointerTy())),
2332 llvm::ConstantInt::get(CGF.Int32Ty, /*V=*/0, /*isSigned=*/true)};
2333 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
2334 CGM.getModule(), OMPRTL___kmpc_barrier_simple_spmd),
2335 Args);
2336}
2337
2338void CGOpenMPRuntimeGPU::emitBarrierCall(CodeGenFunction &CGF,
2339 SourceLocation Loc,
2340 OpenMPDirectiveKind Kind, bool,
2341 bool) {
2342 // Always emit simple barriers!
2343 if (!CGF.HaveInsertPoint())
2344 return;
2345 // Build call __kmpc_cancel_barrier(loc, thread_id);
2346 unsigned Flags = getDefaultFlagsForBarriers(Kind);
2347 llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc, Flags),
2348 getThreadID(CGF, Loc)};
2349
2350 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
2351 CGM.getModule(), OMPRTL___kmpc_barrier),
2352 Args);
2353}
2354
2355void CGOpenMPRuntimeGPU::emitCriticalRegion(
2356 CodeGenFunction &CGF, StringRef CriticalName,
2357 const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc,
2358 const Expr *Hint) {
2359 llvm::BasicBlock *LoopBB = CGF.createBasicBlock("omp.critical.loop");
2360 llvm::BasicBlock *TestBB = CGF.createBasicBlock("omp.critical.test");
2361 llvm::BasicBlock *SyncBB = CGF.createBasicBlock("omp.critical.sync");
2362 llvm::BasicBlock *BodyBB = CGF.createBasicBlock("omp.critical.body");
2363 llvm::BasicBlock *ExitBB = CGF.createBasicBlock("omp.critical.exit");
2364
2365 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
2366
2367 // Get the mask of active threads in the warp.
2368 llvm::Value *Mask = CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
2369 CGM.getModule(), OMPRTL___kmpc_warp_active_thread_mask));
2370 // Fetch team-local id of the thread.
2371 llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
2372
2373 // Get the width of the team.
2374 llvm::Value *TeamWidth = RT.getGPUNumThreads(CGF);
2375
2376 // Initialize the counter variable for the loop.
2377 QualType Int32Ty =
2378 CGF.getContext().getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/0);
2379 Address Counter = CGF.CreateMemTemp(Int32Ty, "critical_counter");
2380 LValue CounterLVal = CGF.MakeAddrLValue(Counter, Int32Ty);
2381 CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.Int32Ty), CounterLVal,
2382 /*isInit=*/true);
2383
2384 // Block checks if loop counter exceeds upper bound.
2385 CGF.EmitBlock(LoopBB);
2386 llvm::Value *CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);
2387 llvm::Value *CmpLoopBound = CGF.Builder.CreateICmpSLT(CounterVal, TeamWidth);
2388 CGF.Builder.CreateCondBr(CmpLoopBound, TestBB, ExitBB);
2389
2390 // Block tests which single thread should execute region, and which threads
2391 // should go straight to synchronisation point.
2392 CGF.EmitBlock(TestBB);
2393 CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);
2394 llvm::Value *CmpThreadToCounter =
2395 CGF.Builder.CreateICmpEQ(ThreadID, CounterVal);
2396 CGF.Builder.CreateCondBr(CmpThreadToCounter, BodyBB, SyncBB);
2397
2398 // Block emits the body of the critical region.
2399 CGF.EmitBlock(BodyBB);
2400
2401 // Output the critical statement.
2402 CGOpenMPRuntime::emitCriticalRegion(CGF, CriticalName, CriticalOpGen, Loc,
2403 Hint);
2404
2405 // After the body surrounded by the critical region, the single executing
2406 // thread will jump to the synchronisation point.
2407 // Block waits for all threads in current team to finish then increments the
2408 // counter variable and returns to the loop.
2409 CGF.EmitBlock(SyncBB);
2410 // Reconverge active threads in the warp.
2411 (void)CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
2412 CGM.getModule(), OMPRTL___kmpc_syncwarp),
2413 Mask);
2414
2415 llvm::Value *IncCounterVal =
2416 CGF.Builder.CreateNSWAdd(CounterVal, CGF.Builder.getInt32(1));
2417 CGF.EmitStoreOfScalar(IncCounterVal, CounterLVal);
2418 CGF.EmitBranch(LoopBB);
2419
2420 // Block that is reached when all threads in the team complete the region.
2421 CGF.EmitBlock(ExitBB, /*IsFinished=*/true);
2422}
2423
2424/// Cast value to the specified type.
2425static llvm::Value *castValueToType(CodeGenFunction &CGF, llvm::Value *Val,
2426 QualType ValTy, QualType CastTy,
2427 SourceLocation Loc) {
2428 assert(!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&((!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&
"Cast type must sized.") ? static_cast<void> (0) : __assert_fail
("!CGF.getContext().getTypeSizeInChars(CastTy).isZero() && \"Cast type must sized.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2429, __PRETTY_FUNCTION__))
2429 "Cast type must sized.")((!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&
"Cast type must sized.") ? static_cast<void> (0) : __assert_fail
("!CGF.getContext().getTypeSizeInChars(CastTy).isZero() && \"Cast type must sized.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2429, __PRETTY_FUNCTION__))
;
2430 assert(!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&((!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&
"Val type must sized.") ? static_cast<void> (0) : __assert_fail
("!CGF.getContext().getTypeSizeInChars(ValTy).isZero() && \"Val type must sized.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2431, __PRETTY_FUNCTION__))
2431 "Val type must sized.")((!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&
"Val type must sized.") ? static_cast<void> (0) : __assert_fail
("!CGF.getContext().getTypeSizeInChars(ValTy).isZero() && \"Val type must sized.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2431, __PRETTY_FUNCTION__))
;
2432 llvm::Type *LLVMCastTy = CGF.ConvertTypeForMem(CastTy);
2433 if (ValTy == CastTy)
2434 return Val;
2435 if (CGF.getContext().getTypeSizeInChars(ValTy) ==
2436 CGF.getContext().getTypeSizeInChars(CastTy))
2437 return CGF.Builder.CreateBitCast(Val, LLVMCastTy);
2438 if (CastTy->isIntegerType() && ValTy->isIntegerType())
2439 return CGF.Builder.CreateIntCast(Val, LLVMCastTy,
2440 CastTy->hasSignedIntegerRepresentation());
2441 Address CastItem = CGF.CreateMemTemp(CastTy);
2442 Address ValCastItem = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
2443 CastItem, Val->getType()->getPointerTo(CastItem.getAddressSpace()));
2444 CGF.EmitStoreOfScalar(Val, ValCastItem, /*Volatile=*/false, ValTy,
2445 LValueBaseInfo(AlignmentSource::Type),
2446 TBAAAccessInfo());
2447 return CGF.EmitLoadOfScalar(CastItem, /*Volatile=*/false, CastTy, Loc,
2448 LValueBaseInfo(AlignmentSource::Type),
2449 TBAAAccessInfo());
2450}
2451
2452/// This function creates calls to one of two shuffle functions to copy
2453/// variables between lanes in a warp.
2454static llvm::Value *createRuntimeShuffleFunction(CodeGenFunction &CGF,
2455 llvm::Value *Elem,
2456 QualType ElemType,
2457 llvm::Value *Offset,
2458 SourceLocation Loc) {
2459 CodeGenModule &CGM = CGF.CGM;
2460 CGBuilderTy &Bld = CGF.Builder;
2461 CGOpenMPRuntimeGPU &RT =
2462 *(static_cast<CGOpenMPRuntimeGPU *>(&CGM.getOpenMPRuntime()));
2463 llvm::OpenMPIRBuilder &OMPBuilder = RT.getOMPBuilder();
2464
2465 CharUnits Size = CGF.getContext().getTypeSizeInChars(ElemType);
2466 assert(Size.getQuantity() <= 8 &&((Size.getQuantity() <= 8 && "Unsupported bitwidth in shuffle instruction."
) ? static_cast<void> (0) : __assert_fail ("Size.getQuantity() <= 8 && \"Unsupported bitwidth in shuffle instruction.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2467, __PRETTY_FUNCTION__))
2467 "Unsupported bitwidth in shuffle instruction.")((Size.getQuantity() <= 8 && "Unsupported bitwidth in shuffle instruction."
) ? static_cast<void> (0) : __assert_fail ("Size.getQuantity() <= 8 && \"Unsupported bitwidth in shuffle instruction.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 2467, __PRETTY_FUNCTION__))
;
2468
2469 RuntimeFunction ShuffleFn = Size.getQuantity() <= 4
2470 ? OMPRTL___kmpc_shuffle_int32
2471 : OMPRTL___kmpc_shuffle_int64;
2472
2473 // Cast all types to 32- or 64-bit values before calling shuffle routines.
2474 QualType CastTy = CGF.getContext().getIntTypeForBitwidth(
2475 Size.getQuantity() <= 4 ? 32 : 64, /*Signed=*/1);
2476 llvm::Value *ElemCast = castValueToType(CGF, Elem, ElemType, CastTy, Loc);
2477 llvm::Value *WarpSize =
2478 Bld.CreateIntCast(RT.getGPUWarpSize(CGF), CGM.Int16Ty, /*isSigned=*/true);
2479
2480 llvm::Value *ShuffledVal = CGF.EmitRuntimeCall(
2481 OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), ShuffleFn),
2482 {ElemCast, Offset, WarpSize});
2483
2484 return castValueToType(CGF, ShuffledVal, CastTy, ElemType, Loc);
2485}
2486
2487static void shuffleAndStore(CodeGenFunction &CGF, Address SrcAddr,
2488 Address DestAddr, QualType ElemType,
2489 llvm::Value *Offset, SourceLocation Loc) {
2490 CGBuilderTy &Bld = CGF.Builder;
2491
2492 CharUnits Size = CGF.getContext().getTypeSizeInChars(ElemType);
2493 // Create the loop over the big sized data.
2494 // ptr = (void*)Elem;
2495 // ptrEnd = (void*) Elem + 1;
2496 // Step = 8;
2497 // while (ptr + Step < ptrEnd)
2498 // shuffle((int64_t)*ptr);
2499 // Step = 4;
2500 // while (ptr + Step < ptrEnd)
2501 // shuffle((int32_t)*ptr);
2502 // ...
2503 Address ElemPtr = DestAddr;
2504 Address Ptr = SrcAddr;
2505 Address PtrEnd = Bld.CreatePointerBitCastOrAddrSpaceCast(
2506 Bld.CreateConstGEP(SrcAddr, 1), CGF.VoidPtrTy);
2507 for (int IntSize = 8; IntSize >= 1; IntSize /= 2) {
2508 if (Size < CharUnits::fromQuantity(IntSize))
2509 continue;
2510 QualType IntType = CGF.getContext().getIntTypeForBitwidth(
2511 CGF.getContext().toBits(CharUnits::fromQuantity(IntSize)),
2512 /*Signed=*/1);
2513 llvm::Type *IntTy = CGF.ConvertTypeForMem(IntType);
2514 Ptr = Bld.CreatePointerBitCastOrAddrSpaceCast(Ptr, IntTy->getPointerTo());
2515 ElemPtr =
2516 Bld.CreatePointerBitCastOrAddrSpaceCast(ElemPtr, IntTy->getPointerTo());
2517 if (Size.getQuantity() / IntSize > 1) {
2518 llvm::BasicBlock *PreCondBB = CGF.createBasicBlock(".shuffle.pre_cond");
2519 llvm::BasicBlock *ThenBB = CGF.createBasicBlock(".shuffle.then");
2520 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".shuffle.exit");
2521 llvm::BasicBlock *CurrentBB = Bld.GetInsertBlock();
2522 CGF.EmitBlock(PreCondBB);
2523 llvm::PHINode *PhiSrc =
2524 Bld.CreatePHI(Ptr.getType(), /*NumReservedValues=*/2);
2525 PhiSrc->addIncoming(Ptr.getPointer(), CurrentBB);
2526 llvm::PHINode *PhiDest =
2527 Bld.CreatePHI(ElemPtr.getType(), /*NumReservedValues=*/2);
2528 PhiDest->addIncoming(ElemPtr.getPointer(), CurrentBB);
2529 Ptr = Address(PhiSrc, Ptr.getAlignment());
2530 ElemPtr = Address(PhiDest, ElemPtr.getAlignment());
2531 llvm::Value *PtrDiff = Bld.CreatePtrDiff(
2532 PtrEnd.getPointer(), Bld.CreatePointerBitCastOrAddrSpaceCast(
2533 Ptr.getPointer(), CGF.VoidPtrTy));
2534 Bld.CreateCondBr(Bld.CreateICmpSGT(PtrDiff, Bld.getInt64(IntSize - 1)),
2535 ThenBB, ExitBB);
2536 CGF.EmitBlock(ThenBB);
2537 llvm::Value *Res = createRuntimeShuffleFunction(
2538 CGF,
2539 CGF.EmitLoadOfScalar(Ptr, /*Volatile=*/false, IntType, Loc,
2540 LValueBaseInfo(AlignmentSource::Type),
2541 TBAAAccessInfo()),
2542 IntType, Offset, Loc);
2543 CGF.EmitStoreOfScalar(Res, ElemPtr, /*Volatile=*/false, IntType,
2544 LValueBaseInfo(AlignmentSource::Type),
2545 TBAAAccessInfo());
2546 Address LocalPtr = Bld.CreateConstGEP(Ptr, 1);
2547 Address LocalElemPtr = Bld.CreateConstGEP(ElemPtr, 1);
2548 PhiSrc->addIncoming(LocalPtr.getPointer(), ThenBB);
2549 PhiDest->addIncoming(LocalElemPtr.getPointer(), ThenBB);
2550 CGF.EmitBranch(PreCondBB);
2551 CGF.EmitBlock(ExitBB);
2552 } else {
2553 llvm::Value *Res = createRuntimeShuffleFunction(
2554 CGF,
2555 CGF.EmitLoadOfScalar(Ptr, /*Volatile=*/false, IntType, Loc,
2556 LValueBaseInfo(AlignmentSource::Type),
2557 TBAAAccessInfo()),
2558 IntType, Offset, Loc);
2559 CGF.EmitStoreOfScalar(Res, ElemPtr, /*Volatile=*/false, IntType,
2560 LValueBaseInfo(AlignmentSource::Type),
2561 TBAAAccessInfo());
2562 Ptr = Bld.CreateConstGEP(Ptr, 1);
2563 ElemPtr = Bld.CreateConstGEP(ElemPtr, 1);
2564 }
2565 Size = Size % IntSize;
2566 }
2567}
2568
2569namespace {
2570enum CopyAction : unsigned {
2571 // RemoteLaneToThread: Copy over a Reduce list from a remote lane in
2572 // the warp using shuffle instructions.
2573 RemoteLaneToThread,
2574 // ThreadCopy: Make a copy of a Reduce list on the thread's stack.
2575 ThreadCopy,
2576 // ThreadToScratchpad: Copy a team-reduced array to the scratchpad.
2577 ThreadToScratchpad,
2578 // ScratchpadToThread: Copy from a scratchpad array in global memory
2579 // containing team-reduced data to a thread's stack.
2580 ScratchpadToThread,
2581};
2582} // namespace
2583
2584struct CopyOptionsTy {
2585 llvm::Value *RemoteLaneOffset;
2586 llvm::Value *ScratchpadIndex;
2587 llvm::Value *ScratchpadWidth;
2588};
2589
2590/// Emit instructions to copy a Reduce list, which contains partially
2591/// aggregated values, in the specified direction.
2592static void emitReductionListCopy(
2593 CopyAction Action, CodeGenFunction &CGF, QualType ReductionArrayTy,
2594 ArrayRef<const Expr *> Privates, Address SrcBase, Address DestBase,
2595 CopyOptionsTy CopyOptions = {nullptr, nullptr, nullptr}) {
2596
2597 CodeGenModule &CGM = CGF.CGM;
2598 ASTContext &C = CGM.getContext();
2599 CGBuilderTy &Bld = CGF.Builder;
2600
2601 llvm::Value *RemoteLaneOffset = CopyOptions.RemoteLaneOffset;
2602 llvm::Value *ScratchpadIndex = CopyOptions.ScratchpadIndex;
2603 llvm::Value *ScratchpadWidth = CopyOptions.ScratchpadWidth;
2604
2605 // Iterates, element-by-element, through the source Reduce list and
2606 // make a copy.
2607 unsigned Idx = 0;
2608 unsigned Size = Privates.size();
2609 for (const Expr *Private : Privates) {
2610 Address SrcElementAddr = Address::invalid();
2611 Address DestElementAddr = Address::invalid();
2612 Address DestElementPtrAddr = Address::invalid();
2613 // Should we shuffle in an element from a remote lane?
2614 bool ShuffleInElement = false;
2615 // Set to true to update the pointer in the dest Reduce list to a
2616 // newly created element.
2617 bool UpdateDestListPtr = false;
2618 // Increment the src or dest pointer to the scratchpad, for each
2619 // new element.
2620 bool IncrScratchpadSrc = false;
2621 bool IncrScratchpadDest = false;
2622
2623 switch (Action) {
2624 case RemoteLaneToThread: {
2625 // Step 1.1: Get the address for the src element in the Reduce list.
2626 Address SrcElementPtrAddr = Bld.CreateConstArrayGEP(SrcBase, Idx);
2627 SrcElementAddr = CGF.EmitLoadOfPointer(
2628 SrcElementPtrAddr,
2629 C.getPointerType(Private->getType())->castAs<PointerType>());
2630
2631 // Step 1.2: Create a temporary to store the element in the destination
2632 // Reduce list.
2633 DestElementPtrAddr = Bld.CreateConstArrayGEP(DestBase, Idx);
2634 DestElementAddr =
2635 CGF.CreateMemTemp(Private->getType(), ".omp.reduction.element");
2636 ShuffleInElement = true;
2637 UpdateDestListPtr = true;
2638 break;
2639 }
2640 case ThreadCopy: {
2641 // Step 1.1: Get the address for the src element in the Reduce list.
2642 Address SrcElementPtrAddr = Bld.CreateConstArrayGEP(SrcBase, Idx);
2643 SrcElementAddr = CGF.EmitLoadOfPointer(
2644 SrcElementPtrAddr,
2645 C.getPointerType(Private->getType())->castAs<PointerType>());
2646
2647 // Step 1.2: Get the address for dest element. The destination
2648 // element has already been created on the thread's stack.
2649 DestElementPtrAddr = Bld.CreateConstArrayGEP(DestBase, Idx);
2650 DestElementAddr = CGF.EmitLoadOfPointer(
2651 DestElementPtrAddr,
2652 C.getPointerType(Private->getType())->castAs<PointerType>());
2653 break;
2654 }
2655 case ThreadToScratchpad: {
2656 // Step 1.1: Get the address for the src element in the Reduce list.
2657 Address SrcElementPtrAddr = Bld.CreateConstArrayGEP(SrcBase, Idx);
2658 SrcElementAddr = CGF.EmitLoadOfPointer(
2659 SrcElementPtrAddr,
2660 C.getPointerType(Private->getType())->castAs<PointerType>());
2661
2662 // Step 1.2: Get the address for dest element:
2663 // address = base + index * ElementSizeInChars.
2664 llvm::Value *ElementSizeInChars = CGF.getTypeSize(Private->getType());
2665 llvm::Value *CurrentOffset =
2666 Bld.CreateNUWMul(ElementSizeInChars, ScratchpadIndex);
2667 llvm::Value *ScratchPadElemAbsolutePtrVal =
2668 Bld.CreateNUWAdd(DestBase.getPointer(), CurrentOffset);
2669 ScratchPadElemAbsolutePtrVal =
2670 Bld.CreateIntToPtr(ScratchPadElemAbsolutePtrVal, CGF.VoidPtrTy);
2671 DestElementAddr = Address(ScratchPadElemAbsolutePtrVal,
2672 C.getTypeAlignInChars(Private->getType()));
2673 IncrScratchpadDest = true;
2674 break;
2675 }
2676 case ScratchpadToThread: {
2677 // Step 1.1: Get the address for the src element in the scratchpad.
2678 // address = base + index * ElementSizeInChars.
2679 llvm::Value *ElementSizeInChars = CGF.getTypeSize(Private->getType());
2680 llvm::Value *CurrentOffset =
2681 Bld.CreateNUWMul(ElementSizeInChars, ScratchpadIndex);
2682 llvm::Value *ScratchPadElemAbsolutePtrVal =
2683 Bld.CreateNUWAdd(SrcBase.getPointer(), CurrentOffset);
2684 ScratchPadElemAbsolutePtrVal =
2685 Bld.CreateIntToPtr(ScratchPadElemAbsolutePtrVal, CGF.VoidPtrTy);
2686 SrcElementAddr = Address(ScratchPadElemAbsolutePtrVal,
2687 C.getTypeAlignInChars(Private->getType()));
2688 IncrScratchpadSrc = true;
2689
2690 // Step 1.2: Create a temporary to store the element in the destination
2691 // Reduce list.
2692 DestElementPtrAddr = Bld.CreateConstArrayGEP(DestBase, Idx);
2693 DestElementAddr =
2694 CGF.CreateMemTemp(Private->getType(), ".omp.reduction.element");
2695 UpdateDestListPtr = true;
2696 break;
2697 }
2698 }
2699
2700 // Regardless of src and dest of copy, we emit the load of src
2701 // element as this is required in all directions
2702 SrcElementAddr = Bld.CreateElementBitCast(
2703 SrcElementAddr, CGF.ConvertTypeForMem(Private->getType()));
2704 DestElementAddr = Bld.CreateElementBitCast(DestElementAddr,
2705 SrcElementAddr.getElementType());
2706
2707 // Now that all active lanes have read the element in the
2708 // Reduce list, shuffle over the value from the remote lane.
2709 if (ShuffleInElement) {
2710 shuffleAndStore(CGF, SrcElementAddr, DestElementAddr, Private->getType(),
2711 RemoteLaneOffset, Private->getExprLoc());
2712 } else {
2713 switch (CGF.getEvaluationKind(Private->getType())) {
2714 case TEK_Scalar: {
2715 llvm::Value *Elem = CGF.EmitLoadOfScalar(
2716 SrcElementAddr, /*Volatile=*/false, Private->getType(),
2717 Private->getExprLoc(), LValueBaseInfo(AlignmentSource::Type),
2718 TBAAAccessInfo());
2719 // Store the source element value to the dest element address.
2720 CGF.EmitStoreOfScalar(
2721 Elem, DestElementAddr, /*Volatile=*/false, Private->getType(),
2722 LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo());
2723 break;
2724 }
2725 case TEK_Complex: {
2726 CodeGenFunction::ComplexPairTy Elem = CGF.EmitLoadOfComplex(
2727 CGF.MakeAddrLValue(SrcElementAddr, Private->getType()),
2728 Private->getExprLoc());
2729 CGF.EmitStoreOfComplex(
2730 Elem, CGF.MakeAddrLValue(DestElementAddr, Private->getType()),
2731 /*isInit=*/false);
2732 break;
2733 }
2734 case TEK_Aggregate:
2735 CGF.EmitAggregateCopy(
2736 CGF.MakeAddrLValue(DestElementAddr, Private->getType()),
2737 CGF.MakeAddrLValue(SrcElementAddr, Private->getType()),
2738 Private->getType(), AggValueSlot::DoesNotOverlap);
2739 break;
2740 }
2741 }
2742
2743 // Step 3.1: Modify reference in dest Reduce list as needed.
2744 // Modifying the reference in Reduce list to point to the newly
2745 // created element. The element is live in the current function
2746 // scope and that of functions it invokes (i.e., reduce_function).
2747 // RemoteReduceData[i] = (void*)&RemoteElem
2748 if (UpdateDestListPtr) {
2749 CGF.EmitStoreOfScalar(Bld.CreatePointerBitCastOrAddrSpaceCast(
2750 DestElementAddr.getPointer(), CGF.VoidPtrTy),
2751 DestElementPtrAddr, /*Volatile=*/false,
2752 C.VoidPtrTy);
2753 }
2754
2755 // Step 4.1: Increment SrcBase/DestBase so that it points to the starting
2756 // address of the next element in scratchpad memory, unless we're currently
2757 // processing the last one. Memory alignment is also taken care of here.
2758 if ((IncrScratchpadDest || IncrScratchpadSrc) && (Idx + 1 < Size)) {
2759 llvm::Value *ScratchpadBasePtr =
2760 IncrScratchpadDest ? DestBase.getPointer() : SrcBase.getPointer();
2761 llvm::Value *ElementSizeInChars = CGF.getTypeSize(Private->getType());
2762 ScratchpadBasePtr = Bld.CreateNUWAdd(
2763 ScratchpadBasePtr,
2764 Bld.CreateNUWMul(ScratchpadWidth, ElementSizeInChars));
2765
2766 // Take care of global memory alignment for performance
2767 ScratchpadBasePtr = Bld.CreateNUWSub(
2768 ScratchpadBasePtr, llvm::ConstantInt::get(CGM.SizeTy, 1));
2769 ScratchpadBasePtr = Bld.CreateUDiv(
2770 ScratchpadBasePtr,
2771 llvm::ConstantInt::get(CGM.SizeTy, GlobalMemoryAlignment));
2772 ScratchpadBasePtr = Bld.CreateNUWAdd(
2773 ScratchpadBasePtr, llvm::ConstantInt::get(CGM.SizeTy, 1));
2774 ScratchpadBasePtr = Bld.CreateNUWMul(
2775 ScratchpadBasePtr,
2776 llvm::ConstantInt::get(CGM.SizeTy, GlobalMemoryAlignment));
2777
2778 if (IncrScratchpadDest)
2779 DestBase = Address(ScratchpadBasePtr, CGF.getPointerAlign());
2780 else /* IncrScratchpadSrc = true */
2781 SrcBase = Address(ScratchpadBasePtr, CGF.getPointerAlign());
2782 }
2783
2784 ++Idx;
2785 }
2786}
2787
2788/// This function emits a helper that gathers Reduce lists from the first
2789/// lane of every active warp to lanes in the first warp.
2790///
2791/// void inter_warp_copy_func(void* reduce_data, num_warps)
2792/// shared smem[warp_size];
2793/// For all data entries D in reduce_data:
2794/// sync
2795/// If (I am the first lane in each warp)
2796/// Copy my local D to smem[warp_id]
2797/// sync
2798/// if (I am the first warp)
2799/// Copy smem[thread_id] to my local D
2800static llvm::Value *emitInterWarpCopyFunction(CodeGenModule &CGM,
2801 ArrayRef<const Expr *> Privates,
2802 QualType ReductionArrayTy,
2803 SourceLocation Loc) {
2804 ASTContext &C = CGM.getContext();
2805 llvm::Module &M = CGM.getModule();
2806
2807 // ReduceList: thread local Reduce list.
2808 // At the stage of the computation when this function is called, partially
2809 // aggregated values reside in the first lane of every active warp.
2810 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2811 C.VoidPtrTy, ImplicitParamDecl::Other);
2812 // NumWarps: number of warps active in the parallel region. This could
2813 // be smaller than 32 (max warps in a CTA) for partial block reduction.
2814 ImplicitParamDecl NumWarpsArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2815 C.getIntTypeForBitwidth(32, /* Signed */ true),
2816 ImplicitParamDecl::Other);
2817 FunctionArgList Args;
2818 Args.push_back(&ReduceListArg);
2819 Args.push_back(&NumWarpsArg);
2820
2821 const CGFunctionInfo &CGFI =
2822 CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
2823 auto *Fn = llvm::Function::Create(CGM.getTypes().GetFunctionType(CGFI),
2824 llvm::GlobalValue::InternalLinkage,
2825 "_omp_reduction_inter_warp_copy_func", &M);
2826 CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
2827 Fn->setDoesNotRecurse();
2828 CodeGenFunction CGF(CGM);
2829 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
2830
2831 CGBuilderTy &Bld = CGF.Builder;
2832
2833 // This array is used as a medium to transfer, one reduce element at a time,
2834 // the data from the first lane of every warp to lanes in the first warp
2835 // in order to perform the final step of a reduction in a parallel region
2836 // (reduction across warps). The array is placed in NVPTX __shared__ memory
2837 // for reduced latency, as well as to have a distinct copy for concurrently
2838 // executing target regions. The array is declared with common linkage so
2839 // as to be shared across compilation units.
2840 StringRef TransferMediumName =
2841 "__openmp_nvptx_data_transfer_temporary_storage";
2842 llvm::GlobalVariable *TransferMedium =
2843 M.getGlobalVariable(TransferMediumName);
2844 unsigned WarpSize = CGF.getTarget().getGridValue(llvm::omp::GV_Warp_Size);
2845 if (!TransferMedium) {
2846 auto *Ty = llvm::ArrayType::get(CGM.Int32Ty, WarpSize);
2847 unsigned SharedAddressSpace = C.getTargetAddressSpace(LangAS::cuda_shared);
2848 TransferMedium = new llvm::GlobalVariable(
2849 M, Ty, /*isConstant=*/false, llvm::GlobalVariable::WeakAnyLinkage,
2850 llvm::UndefValue::get(Ty), TransferMediumName,
2851 /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
2852 SharedAddressSpace);
2853 CGM.addCompilerUsedGlobal(TransferMedium);
2854 }
2855
2856 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
2857 // Get the CUDA thread id of the current OpenMP thread on the GPU.
2858 llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
2859 // nvptx_lane_id = nvptx_id % warpsize
2860 llvm::Value *LaneID = getNVPTXLaneID(CGF);
2861 // nvptx_warp_id = nvptx_id / warpsize
2862 llvm::Value *WarpID = getNVPTXWarpID(CGF);
2863
2864 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
2865 Address LocalReduceList(
2866 Bld.CreatePointerBitCastOrAddrSpaceCast(
2867 CGF.EmitLoadOfScalar(
2868 AddrReduceListArg, /*Volatile=*/false, C.VoidPtrTy, Loc,
2869 LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo()),
2870 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
2871 CGF.getPointerAlign());
2872
2873 unsigned Idx = 0;
2874 for (const Expr *Private : Privates) {
2875 //
2876 // Warp master copies reduce element to transfer medium in __shared__
2877 // memory.
2878 //
2879 unsigned RealTySize =
2880 C.getTypeSizeInChars(Private->getType())
2881 .alignTo(C.getTypeAlignInChars(Private->getType()))
2882 .getQuantity();
2883 for (unsigned TySize = 4; TySize > 0 && RealTySize > 0; TySize /=2) {
2884 unsigned NumIters = RealTySize / TySize;
2885 if (NumIters == 0)
2886 continue;
2887 QualType CType = C.getIntTypeForBitwidth(
2888 C.toBits(CharUnits::fromQuantity(TySize)), /*Signed=*/1);
2889 llvm::Type *CopyType = CGF.ConvertTypeForMem(CType);
2890 CharUnits Align = CharUnits::fromQuantity(TySize);
2891 llvm::Value *Cnt = nullptr;
2892 Address CntAddr = Address::invalid();
2893 llvm::BasicBlock *PrecondBB = nullptr;
2894 llvm::BasicBlock *ExitBB = nullptr;
2895 if (NumIters > 1) {
2896 CntAddr = CGF.CreateMemTemp(C.IntTy, ".cnt.addr");
2897 CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.IntTy), CntAddr,
2898 /*Volatile=*/false, C.IntTy);
2899 PrecondBB = CGF.createBasicBlock("precond");
2900 ExitBB = CGF.createBasicBlock("exit");
2901 llvm::BasicBlock *BodyBB = CGF.createBasicBlock("body");
2902 // There is no need to emit line number for unconditional branch.
2903 (void)ApplyDebugLocation::CreateEmpty(CGF);
2904 CGF.EmitBlock(PrecondBB);
2905 Cnt = CGF.EmitLoadOfScalar(CntAddr, /*Volatile=*/false, C.IntTy, Loc);
2906 llvm::Value *Cmp =
2907 Bld.CreateICmpULT(Cnt, llvm::ConstantInt::get(CGM.IntTy, NumIters));
2908 Bld.CreateCondBr(Cmp, BodyBB, ExitBB);
2909 CGF.EmitBlock(BodyBB);
2910 }
2911 // kmpc_barrier.
2912 CGM.getOpenMPRuntime().emitBarrierCall(CGF, Loc, OMPD_unknown,
2913 /*EmitChecks=*/false,
2914 /*ForceSimpleCall=*/true);
2915 llvm::BasicBlock *ThenBB = CGF.createBasicBlock("then");
2916 llvm::BasicBlock *ElseBB = CGF.createBasicBlock("else");
2917 llvm::BasicBlock *MergeBB = CGF.createBasicBlock("ifcont");
2918
2919 // if (lane_id == 0)
2920 llvm::Value *IsWarpMaster = Bld.CreateIsNull(LaneID, "warp_master");
2921 Bld.CreateCondBr(IsWarpMaster, ThenBB, ElseBB);
2922 CGF.EmitBlock(ThenBB);
2923
2924 // Reduce element = LocalReduceList[i]
2925 Address ElemPtrPtrAddr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
2926 llvm::Value *ElemPtrPtr = CGF.EmitLoadOfScalar(
2927 ElemPtrPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
2928 // elemptr = ((CopyType*)(elemptrptr)) + I
2929 Address ElemPtr = Address(ElemPtrPtr, Align);
2930 ElemPtr = Bld.CreateElementBitCast(ElemPtr, CopyType);
2931 if (NumIters > 1) {
2932 ElemPtr = Address(Bld.CreateGEP(ElemPtr.getPointer(), Cnt),
2933 ElemPtr.getAlignment());
2934 }
2935
2936 // Get pointer to location in transfer medium.
2937 // MediumPtr = &medium[warp_id]
2938 llvm::Value *MediumPtrVal = Bld.CreateInBoundsGEP(
2939 TransferMedium->getValueType(), TransferMedium,
2940 {llvm::Constant::getNullValue(CGM.Int64Ty), WarpID});
2941 Address MediumPtr(MediumPtrVal, Align);
2942 // Casting to actual data type.
2943 // MediumPtr = (CopyType*)MediumPtrAddr;
2944 MediumPtr = Bld.CreateElementBitCast(MediumPtr, CopyType);
2945
2946 // elem = *elemptr
2947 //*MediumPtr = elem
2948 llvm::Value *Elem = CGF.EmitLoadOfScalar(
2949 ElemPtr, /*Volatile=*/false, CType, Loc,
2950 LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo());
2951 // Store the source element value to the dest element address.
2952 CGF.EmitStoreOfScalar(Elem, MediumPtr, /*Volatile=*/true, CType,
2953 LValueBaseInfo(AlignmentSource::Type),
2954 TBAAAccessInfo());
2955
2956 Bld.CreateBr(MergeBB);
2957
2958 CGF.EmitBlock(ElseBB);
2959 Bld.CreateBr(MergeBB);
2960
2961 CGF.EmitBlock(MergeBB);
2962
2963 // kmpc_barrier.
2964 CGM.getOpenMPRuntime().emitBarrierCall(CGF, Loc, OMPD_unknown,
2965 /*EmitChecks=*/false,
2966 /*ForceSimpleCall=*/true);
2967
2968 //
2969 // Warp 0 copies reduce element from transfer medium.
2970 //
2971 llvm::BasicBlock *W0ThenBB = CGF.createBasicBlock("then");
2972 llvm::BasicBlock *W0ElseBB = CGF.createBasicBlock("else");
2973 llvm::BasicBlock *W0MergeBB = CGF.createBasicBlock("ifcont");
2974
2975 Address AddrNumWarpsArg = CGF.GetAddrOfLocalVar(&NumWarpsArg);
2976 llvm::Value *NumWarpsVal = CGF.EmitLoadOfScalar(
2977 AddrNumWarpsArg, /*Volatile=*/false, C.IntTy, Loc);
2978
2979 // Up to 32 threads in warp 0 are active.
2980 llvm::Value *IsActiveThread =
2981 Bld.CreateICmpULT(ThreadID, NumWarpsVal, "is_active_thread");
2982 Bld.CreateCondBr(IsActiveThread, W0ThenBB, W0ElseBB);
2983
2984 CGF.EmitBlock(W0ThenBB);
2985
2986 // SrcMediumPtr = &medium[tid]
2987 llvm::Value *SrcMediumPtrVal = Bld.CreateInBoundsGEP(
2988 TransferMedium->getValueType(), TransferMedium,
2989 {llvm::Constant::getNullValue(CGM.Int64Ty), ThreadID});
2990 Address SrcMediumPtr(SrcMediumPtrVal, Align);
2991 // SrcMediumVal = *SrcMediumPtr;
2992 SrcMediumPtr = Bld.CreateElementBitCast(SrcMediumPtr, CopyType);
2993
2994 // TargetElemPtr = (CopyType*)(SrcDataAddr[i]) + I
2995 Address TargetElemPtrPtr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
2996 llvm::Value *TargetElemPtrVal = CGF.EmitLoadOfScalar(
2997 TargetElemPtrPtr, /*Volatile=*/false, C.VoidPtrTy, Loc);
2998 Address TargetElemPtr = Address(TargetElemPtrVal, Align);
2999 TargetElemPtr = Bld.CreateElementBitCast(TargetElemPtr, CopyType);
3000 if (NumIters > 1) {
3001 TargetElemPtr = Address(Bld.CreateGEP(TargetElemPtr.getPointer(), Cnt),
3002 TargetElemPtr.getAlignment());
3003 }
3004
3005 // *TargetElemPtr = SrcMediumVal;
3006 llvm::Value *SrcMediumValue =
3007 CGF.EmitLoadOfScalar(SrcMediumPtr, /*Volatile=*/true, CType, Loc);
3008 CGF.EmitStoreOfScalar(SrcMediumValue, TargetElemPtr, /*Volatile=*/false,
3009 CType);
3010 Bld.CreateBr(W0MergeBB);
3011
3012 CGF.EmitBlock(W0ElseBB);
3013 Bld.CreateBr(W0MergeBB);
3014
3015 CGF.EmitBlock(W0MergeBB);
3016
3017 if (NumIters > 1) {
3018 Cnt = Bld.CreateNSWAdd(Cnt, llvm::ConstantInt::get(CGM.IntTy, /*V=*/1));
3019 CGF.EmitStoreOfScalar(Cnt, CntAddr, /*Volatile=*/false, C.IntTy);
3020 CGF.EmitBranch(PrecondBB);
3021 (void)ApplyDebugLocation::CreateEmpty(CGF);
3022 CGF.EmitBlock(ExitBB);
3023 }
3024 RealTySize %= TySize;
3025 }
3026 ++Idx;
3027 }
3028
3029 CGF.FinishFunction();
3030 return Fn;
3031}
3032
3033/// Emit a helper that reduces data across two OpenMP threads (lanes)
3034/// in the same warp. It uses shuffle instructions to copy over data from
3035/// a remote lane's stack. The reduction algorithm performed is specified
3036/// by the fourth parameter.
3037///
3038/// Algorithm Versions.
3039/// Full Warp Reduce (argument value 0):
3040/// This algorithm assumes that all 32 lanes are active and gathers
3041/// data from these 32 lanes, producing a single resultant value.
3042/// Contiguous Partial Warp Reduce (argument value 1):
3043/// This algorithm assumes that only a *contiguous* subset of lanes
3044/// are active. This happens for the last warp in a parallel region
3045/// when the user specified num_threads is not an integer multiple of
3046/// 32. This contiguous subset always starts with the zeroth lane.
3047/// Partial Warp Reduce (argument value 2):
3048/// This algorithm gathers data from any number of lanes at any position.
3049/// All reduced values are stored in the lowest possible lane. The set
3050/// of problems every algorithm addresses is a super set of those
3051/// addressable by algorithms with a lower version number. Overhead
3052/// increases as algorithm version increases.
3053///
3054/// Terminology
3055/// Reduce element:
3056/// Reduce element refers to the individual data field with primitive
3057/// data types to be combined and reduced across threads.
3058/// Reduce list:
3059/// Reduce list refers to a collection of local, thread-private
3060/// reduce elements.
3061/// Remote Reduce list:
3062/// Remote Reduce list refers to a collection of remote (relative to
3063/// the current thread) reduce elements.
3064///
3065/// We distinguish between three states of threads that are important to
3066/// the implementation of this function.
3067/// Alive threads:
3068/// Threads in a warp executing the SIMT instruction, as distinguished from
3069/// threads that are inactive due to divergent control flow.
3070/// Active threads:
3071/// The minimal set of threads that has to be alive upon entry to this
3072/// function. The computation is correct iff active threads are alive.
3073/// Some threads are alive but they are not active because they do not
3074/// contribute to the computation in any useful manner. Turning them off
3075/// may introduce control flow overheads without any tangible benefits.
3076/// Effective threads:
3077/// In order to comply with the argument requirements of the shuffle
3078/// function, we must keep all lanes holding data alive. But at most
3079/// half of them perform value aggregation; we refer to this half of
3080/// threads as effective. The other half is simply handing off their
3081/// data.
3082///
3083/// Procedure
3084/// Value shuffle:
3085/// In this step active threads transfer data from higher lane positions
3086/// in the warp to lower lane positions, creating Remote Reduce list.
3087/// Value aggregation:
3088/// In this step, effective threads combine their thread local Reduce list
3089/// with Remote Reduce list and store the result in the thread local
3090/// Reduce list.
3091/// Value copy:
3092/// In this step, we deal with the assumption made by algorithm 2
3093/// (i.e. contiguity assumption). When we have an odd number of lanes
3094/// active, say 2k+1, only k threads will be effective and therefore k
3095/// new values will be produced. However, the Reduce list owned by the
3096/// (2k+1)th thread is ignored in the value aggregation. Therefore
3097/// we copy the Reduce list from the (2k+1)th lane to (k+1)th lane so
3098/// that the contiguity assumption still holds.
3099static llvm::Function *emitShuffleAndReduceFunction(
3100 CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
3101 QualType ReductionArrayTy, llvm::Function *ReduceFn, SourceLocation Loc) {
3102 ASTContext &C = CGM.getContext();
3103
3104 // Thread local Reduce list used to host the values of data to be reduced.
3105 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3106 C.VoidPtrTy, ImplicitParamDecl::Other);
3107 // Current lane id; could be logical.
3108 ImplicitParamDecl LaneIDArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.ShortTy,
3109 ImplicitParamDecl::Other);
3110 // Offset of the remote source lane relative to the current lane.
3111 ImplicitParamDecl RemoteLaneOffsetArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3112 C.ShortTy, ImplicitParamDecl::Other);
3113 // Algorithm version. This is expected to be known at compile time.
3114 ImplicitParamDecl AlgoVerArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3115 C.ShortTy, ImplicitParamDecl::Other);
3116 FunctionArgList Args;
3117 Args.push_back(&ReduceListArg);
3118 Args.push_back(&LaneIDArg);
3119 Args.push_back(&RemoteLaneOffsetArg);
3120 Args.push_back(&AlgoVerArg);
3121
3122 const CGFunctionInfo &CGFI =
3123 CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
3124 auto *Fn = llvm::Function::Create(
3125 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
3126 "_omp_reduction_shuffle_and_reduce_func", &CGM.getModule());
3127 CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
3128 Fn->setDoesNotRecurse();
3129
3130 CodeGenFunction CGF(CGM);
3131 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
3132
3133 CGBuilderTy &Bld = CGF.Builder;
3134
3135 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
3136 Address LocalReduceList(
3137 Bld.CreatePointerBitCastOrAddrSpaceCast(
3138 CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
3139 C.VoidPtrTy, SourceLocation()),
3140 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
3141 CGF.getPointerAlign());
3142
3143 Address AddrLaneIDArg = CGF.GetAddrOfLocalVar(&LaneIDArg);
3144 llvm::Value *LaneIDArgVal = CGF.EmitLoadOfScalar(
3145 AddrLaneIDArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
3146
3147 Address AddrRemoteLaneOffsetArg = CGF.GetAddrOfLocalVar(&RemoteLaneOffsetArg);
3148 llvm::Value *RemoteLaneOffsetArgVal = CGF.EmitLoadOfScalar(
3149 AddrRemoteLaneOffsetArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
3150
3151 Address AddrAlgoVerArg = CGF.GetAddrOfLocalVar(&AlgoVerArg);
3152 llvm::Value *AlgoVerArgVal = CGF.EmitLoadOfScalar(
3153 AddrAlgoVerArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
3154
3155 // Create a local thread-private variable to host the Reduce list
3156 // from a remote lane.
3157 Address RemoteReduceList =
3158 CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.remote_reduce_list");
3159
3160 // This loop iterates through the list of reduce elements and copies,
3161 // element by element, from a remote lane in the warp to RemoteReduceList,
3162 // hosted on the thread's stack.
3163 emitReductionListCopy(RemoteLaneToThread, CGF, ReductionArrayTy, Privates,
3164 LocalReduceList, RemoteReduceList,
3165 {/*RemoteLaneOffset=*/RemoteLaneOffsetArgVal,
3166 /*ScratchpadIndex=*/nullptr,
3167 /*ScratchpadWidth=*/nullptr});
3168
3169 // The actions to be performed on the Remote Reduce list is dependent
3170 // on the algorithm version.
3171 //
3172 // if (AlgoVer==0) || (AlgoVer==1 && (LaneId < Offset)) || (AlgoVer==2 &&
3173 // LaneId % 2 == 0 && Offset > 0):
3174 // do the reduction value aggregation
3175 //
3176 // The thread local variable Reduce list is mutated in place to host the
3177 // reduced data, which is the aggregated value produced from local and
3178 // remote lanes.
3179 //
3180 // Note that AlgoVer is expected to be a constant integer known at compile
3181 // time.
3182 // When AlgoVer==0, the first conjunction evaluates to true, making
3183 // the entire predicate true during compile time.
3184 // When AlgoVer==1, the second conjunction has only the second part to be
3185 // evaluated during runtime. Other conjunctions evaluates to false
3186 // during compile time.
3187 // When AlgoVer==2, the third conjunction has only the second part to be
3188 // evaluated during runtime. Other conjunctions evaluates to false
3189 // during compile time.
3190 llvm::Value *CondAlgo0 = Bld.CreateIsNull(AlgoVerArgVal);
3191
3192 llvm::Value *Algo1 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(1));
3193 llvm::Value *CondAlgo1 = Bld.CreateAnd(
3194 Algo1, Bld.CreateICmpULT(LaneIDArgVal, RemoteLaneOffsetArgVal));
3195
3196 llvm::Value *Algo2 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(2));
3197 llvm::Value *CondAlgo2 = Bld.CreateAnd(
3198 Algo2, Bld.CreateIsNull(Bld.CreateAnd(LaneIDArgVal, Bld.getInt16(1))));
3199 CondAlgo2 = Bld.CreateAnd(
3200 CondAlgo2, Bld.CreateICmpSGT(RemoteLaneOffsetArgVal, Bld.getInt16(0)));
3201
3202 llvm::Value *CondReduce = Bld.CreateOr(CondAlgo0, CondAlgo1);
3203 CondReduce = Bld.CreateOr(CondReduce, CondAlgo2);
3204
3205 llvm::BasicBlock *ThenBB = CGF.createBasicBlock("then");
3206 llvm::BasicBlock *ElseBB = CGF.createBasicBlock("else");
3207 llvm::BasicBlock *MergeBB = CGF.createBasicBlock("ifcont");
3208 Bld.CreateCondBr(CondReduce, ThenBB, ElseBB);
3209
3210 CGF.EmitBlock(ThenBB);
3211 // reduce_function(LocalReduceList, RemoteReduceList)
3212 llvm::Value *LocalReduceListPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3213 LocalReduceList.getPointer(), CGF.VoidPtrTy);
3214 llvm::Value *RemoteReduceListPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3215 RemoteReduceList.getPointer(), CGF.VoidPtrTy);
3216 CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
3217 CGF, Loc, ReduceFn, {LocalReduceListPtr, RemoteReduceListPtr});
3218 Bld.CreateBr(MergeBB);
3219
3220 CGF.EmitBlock(ElseBB);
3221 Bld.CreateBr(MergeBB);
3222
3223 CGF.EmitBlock(MergeBB);
3224
3225 // if (AlgoVer==1 && (LaneId >= Offset)) copy Remote Reduce list to local
3226 // Reduce list.
3227 Algo1 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(1));
3228 llvm::Value *CondCopy = Bld.CreateAnd(
3229 Algo1, Bld.CreateICmpUGE(LaneIDArgVal, RemoteLaneOffsetArgVal));
3230
3231 llvm::BasicBlock *CpyThenBB = CGF.createBasicBlock("then");
3232 llvm::BasicBlock *CpyElseBB = CGF.createBasicBlock("else");
3233 llvm::BasicBlock *CpyMergeBB = CGF.createBasicBlock("ifcont");
3234 Bld.CreateCondBr(CondCopy, CpyThenBB, CpyElseBB);
3235
3236 CGF.EmitBlock(CpyThenBB);
3237 emitReductionListCopy(ThreadCopy, CGF, ReductionArrayTy, Privates,
3238 RemoteReduceList, LocalReduceList);
3239 Bld.CreateBr(CpyMergeBB);
3240
3241 CGF.EmitBlock(CpyElseBB);
3242 Bld.CreateBr(CpyMergeBB);
3243
3244 CGF.EmitBlock(CpyMergeBB);
3245
3246 CGF.FinishFunction();
3247 return Fn;
3248}
3249
3250/// This function emits a helper that copies all the reduction variables from
3251/// the team into the provided global buffer for the reduction variables.
3252///
3253/// void list_to_global_copy_func(void *buffer, int Idx, void *reduce_data)
3254/// For all data entries D in reduce_data:
3255/// Copy local D to buffer.D[Idx]
3256static llvm::Value *emitListToGlobalCopyFunction(
3257 CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
3258 QualType ReductionArrayTy, SourceLocation Loc,
3259 const RecordDecl *TeamReductionRec,
3260 const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
3261 &VarFieldMap) {
3262 ASTContext &C = CGM.getContext();
3263
3264 // Buffer: global reduction buffer.
3265 ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3266 C.VoidPtrTy, ImplicitParamDecl::Other);
3267 // Idx: index of the buffer.
3268 ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
3269 ImplicitParamDecl::Other);
3270 // ReduceList: thread local Reduce list.
3271 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3272 C.VoidPtrTy, ImplicitParamDecl::Other);
3273 FunctionArgList Args;
3274 Args.push_back(&BufferArg);
3275 Args.push_back(&IdxArg);
3276 Args.push_back(&ReduceListArg);
3277
3278 const CGFunctionInfo &CGFI =
3279 CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
3280 auto *Fn = llvm::Function::Create(
3281 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
3282 "_omp_reduction_list_to_global_copy_func", &CGM.getModule());
3283 CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
3284 Fn->setDoesNotRecurse();
3285 CodeGenFunction CGF(CGM);
3286 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
3287
3288 CGBuilderTy &Bld = CGF.Builder;
3289
3290 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
3291 Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
3292 Address LocalReduceList(
3293 Bld.CreatePointerBitCastOrAddrSpaceCast(
3294 CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
3295 C.VoidPtrTy, Loc),
3296 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
3297 CGF.getPointerAlign());
3298 QualType StaticTy = C.getRecordType(TeamReductionRec);
3299 llvm::Type *LLVMReductionsBufferTy =
3300 CGM.getTypes().ConvertTypeForMem(StaticTy);
3301 llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3302 CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
3303 LLVMReductionsBufferTy->getPointerTo());
3304 llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
3305 CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
3306 /*Volatile=*/false, C.IntTy,
3307 Loc)};
3308 unsigned Idx = 0;
3309 for (const Expr *Private : Privates) {
3310 // Reduce element = LocalReduceList[i]
3311 Address ElemPtrPtrAddr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
3312 llvm::Value *ElemPtrPtr = CGF.EmitLoadOfScalar(
3313 ElemPtrPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
3314 // elemptr = ((CopyType*)(elemptrptr)) + I
3315 ElemPtrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3316 ElemPtrPtr, CGF.ConvertTypeForMem(Private->getType())->getPointerTo());
3317 Address ElemPtr =
3318 Address(ElemPtrPtr, C.getTypeAlignInChars(Private->getType()));
3319 const ValueDecl *VD = cast<DeclRefExpr>(Private)->getDecl();
3320 // Global = Buffer.VD[Idx];
3321 const FieldDecl *FD = VarFieldMap.lookup(VD);
3322 LValue GlobLVal = CGF.EmitLValueForField(
3323 CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
3324 Address GlobAddr = GlobLVal.getAddress(CGF);
3325 llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
3326 GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
3327 GlobLVal.setAddress(Address(BufferPtr, GlobAddr.getAlignment()));
3328 switch (CGF.getEvaluationKind(Private->getType())) {
3329 case TEK_Scalar: {
3330 llvm::Value *V = CGF.EmitLoadOfScalar(
3331 ElemPtr, /*Volatile=*/false, Private->getType(), Loc,
3332 LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo());
3333 CGF.EmitStoreOfScalar(V, GlobLVal);
3334 break;
3335 }
3336 case TEK_Complex: {
3337 CodeGenFunction::ComplexPairTy V = CGF.EmitLoadOfComplex(
3338 CGF.MakeAddrLValue(ElemPtr, Private->getType()), Loc);
3339 CGF.EmitStoreOfComplex(V, GlobLVal, /*isInit=*/false);
3340 break;
3341 }
3342 case TEK_Aggregate:
3343 CGF.EmitAggregateCopy(GlobLVal,
3344 CGF.MakeAddrLValue(ElemPtr, Private->getType()),
3345 Private->getType(), AggValueSlot::DoesNotOverlap);
3346 break;
3347 }
3348 ++Idx;
3349 }
3350
3351 CGF.FinishFunction();
3352 return Fn;
3353}
3354
3355/// This function emits a helper that reduces all the reduction variables from
3356/// the team into the provided global buffer for the reduction variables.
3357///
3358/// void list_to_global_reduce_func(void *buffer, int Idx, void *reduce_data)
3359/// void *GlobPtrs[];
3360/// GlobPtrs[0] = (void*)&buffer.D0[Idx];
3361/// ...
3362/// GlobPtrs[N] = (void*)&buffer.DN[Idx];
3363/// reduce_function(GlobPtrs, reduce_data);
3364static llvm::Value *emitListToGlobalReduceFunction(
3365 CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
3366 QualType ReductionArrayTy, SourceLocation Loc,
3367 const RecordDecl *TeamReductionRec,
3368 const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
3369 &VarFieldMap,
3370 llvm::Function *ReduceFn) {
3371 ASTContext &C = CGM.getContext();
3372
3373 // Buffer: global reduction buffer.
3374 ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3375 C.VoidPtrTy, ImplicitParamDecl::Other);
3376 // Idx: index of the buffer.
3377 ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
3378 ImplicitParamDecl::Other);
3379 // ReduceList: thread local Reduce list.
3380 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3381 C.VoidPtrTy, ImplicitParamDecl::Other);
3382 FunctionArgList Args;
3383 Args.push_back(&BufferArg);
3384 Args.push_back(&IdxArg);
3385 Args.push_back(&ReduceListArg);
3386
3387 const CGFunctionInfo &CGFI =
3388 CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
3389 auto *Fn = llvm::Function::Create(
3390 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
3391 "_omp_reduction_list_to_global_reduce_func", &CGM.getModule());
3392 CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
3393 Fn->setDoesNotRecurse();
3394 CodeGenFunction CGF(CGM);
3395 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
3396
3397 CGBuilderTy &Bld = CGF.Builder;
3398
3399 Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
3400 QualType StaticTy = C.getRecordType(TeamReductionRec);
3401 llvm::Type *LLVMReductionsBufferTy =
3402 CGM.getTypes().ConvertTypeForMem(StaticTy);
3403 llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3404 CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
3405 LLVMReductionsBufferTy->getPointerTo());
3406
3407 // 1. Build a list of reduction variables.
3408 // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
3409 Address ReductionList =
3410 CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.red_list");
3411 auto IPriv = Privates.begin();
3412 llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
3413 CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
3414 /*Volatile=*/false, C.IntTy,
3415 Loc)};
3416 unsigned Idx = 0;
3417 for (unsigned I = 0, E = Privates.size(); I < E; ++I, ++IPriv, ++Idx) {
3418 Address Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3419 // Global = Buffer.VD[Idx];
3420 const ValueDecl *VD = cast<DeclRefExpr>(*IPriv)->getDecl();
3421 const FieldDecl *FD = VarFieldMap.lookup(VD);
3422 LValue GlobLVal = CGF.EmitLValueForField(
3423 CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
3424 Address GlobAddr = GlobLVal.getAddress(CGF);
3425 llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
3426 GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
3427 llvm::Value *Ptr = CGF.EmitCastToVoidPtr(BufferPtr);
3428 CGF.EmitStoreOfScalar(Ptr, Elem, /*Volatile=*/false, C.VoidPtrTy);
3429 if ((*IPriv)->getType()->isVariablyModifiedType()) {
3430 // Store array size.
3431 ++Idx;
3432 Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3433 llvm::Value *Size = CGF.Builder.CreateIntCast(
3434 CGF.getVLASize(
3435 CGF.getContext().getAsVariableArrayType((*IPriv)->getType()))
3436 .NumElts,
3437 CGF.SizeTy, /*isSigned=*/false);
3438 CGF.Builder.CreateStore(CGF.Builder.CreateIntToPtr(Size, CGF.VoidPtrTy),
3439 Elem);
3440 }
3441 }
3442
3443 // Call reduce_function(GlobalReduceList, ReduceList)
3444 llvm::Value *GlobalReduceList =
3445 CGF.EmitCastToVoidPtr(ReductionList.getPointer());
3446 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
3447 llvm::Value *ReducedPtr = CGF.EmitLoadOfScalar(
3448 AddrReduceListArg, /*Volatile=*/false, C.VoidPtrTy, Loc);
3449 CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
3450 CGF, Loc, ReduceFn, {GlobalReduceList, ReducedPtr});
3451 CGF.FinishFunction();
3452 return Fn;
3453}
3454
3455/// This function emits a helper that copies all the reduction variables from
3456/// the team into the provided global buffer for the reduction variables.
3457///
3458/// void list_to_global_copy_func(void *buffer, int Idx, void *reduce_data)
3459/// For all data entries D in reduce_data:
3460/// Copy buffer.D[Idx] to local D;
3461static llvm::Value *emitGlobalToListCopyFunction(
3462 CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
3463 QualType ReductionArrayTy, SourceLocation Loc,
3464 const RecordDecl *TeamReductionRec,
3465 const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
3466 &VarFieldMap) {
3467 ASTContext &C = CGM.getContext();
3468
3469 // Buffer: global reduction buffer.
3470 ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3471 C.VoidPtrTy, ImplicitParamDecl::Other);
3472 // Idx: index of the buffer.
3473 ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
3474 ImplicitParamDecl::Other);
3475 // ReduceList: thread local Reduce list.
3476 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3477 C.VoidPtrTy, ImplicitParamDecl::Other);
3478 FunctionArgList Args;
3479 Args.push_back(&BufferArg);
3480 Args.push_back(&IdxArg);
3481 Args.push_back(&ReduceListArg);
3482
3483 const CGFunctionInfo &CGFI =
3484 CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
3485 auto *Fn = llvm::Function::Create(
3486 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
3487 "_omp_reduction_global_to_list_copy_func", &CGM.getModule());
3488 CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
3489 Fn->setDoesNotRecurse();
3490 CodeGenFunction CGF(CGM);
3491 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
3492
3493 CGBuilderTy &Bld = CGF.Builder;
3494
3495 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
3496 Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
3497 Address LocalReduceList(
3498 Bld.CreatePointerBitCastOrAddrSpaceCast(
3499 CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
3500 C.VoidPtrTy, Loc),
3501 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
3502 CGF.getPointerAlign());
3503 QualType StaticTy = C.getRecordType(TeamReductionRec);
3504 llvm::Type *LLVMReductionsBufferTy =
3505 CGM.getTypes().ConvertTypeForMem(StaticTy);
3506 llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3507 CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
3508 LLVMReductionsBufferTy->getPointerTo());
3509
3510 llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
3511 CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
3512 /*Volatile=*/false, C.IntTy,
3513 Loc)};
3514 unsigned Idx = 0;
3515 for (const Expr *Private : Privates) {
3516 // Reduce element = LocalReduceList[i]
3517 Address ElemPtrPtrAddr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
3518 llvm::Value *ElemPtrPtr = CGF.EmitLoadOfScalar(
3519 ElemPtrPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
3520 // elemptr = ((CopyType*)(elemptrptr)) + I
3521 ElemPtrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3522 ElemPtrPtr, CGF.ConvertTypeForMem(Private->getType())->getPointerTo());
3523 Address ElemPtr =
3524 Address(ElemPtrPtr, C.getTypeAlignInChars(Private->getType()));
3525 const ValueDecl *VD = cast<DeclRefExpr>(Private)->getDecl();
3526 // Global = Buffer.VD[Idx];
3527 const FieldDecl *FD = VarFieldMap.lookup(VD);
3528 LValue GlobLVal = CGF.EmitLValueForField(
3529 CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
3530 Address GlobAddr = GlobLVal.getAddress(CGF);
3531 llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
3532 GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
3533 GlobLVal.setAddress(Address(BufferPtr, GlobAddr.getAlignment()));
3534 switch (CGF.getEvaluationKind(Private->getType())) {
3535 case TEK_Scalar: {
3536 llvm::Value *V = CGF.EmitLoadOfScalar(GlobLVal, Loc);
3537 CGF.EmitStoreOfScalar(V, ElemPtr, /*Volatile=*/false, Private->getType(),
3538 LValueBaseInfo(AlignmentSource::Type),
3539 TBAAAccessInfo());
3540 break;
3541 }
3542 case TEK_Complex: {
3543 CodeGenFunction::ComplexPairTy V = CGF.EmitLoadOfComplex(GlobLVal, Loc);
3544 CGF.EmitStoreOfComplex(V, CGF.MakeAddrLValue(ElemPtr, Private->getType()),
3545 /*isInit=*/false);
3546 break;
3547 }
3548 case TEK_Aggregate:
3549 CGF.EmitAggregateCopy(CGF.MakeAddrLValue(ElemPtr, Private->getType()),
3550 GlobLVal, Private->getType(),
3551 AggValueSlot::DoesNotOverlap);
3552 break;
3553 }
3554 ++Idx;
3555 }
3556
3557 CGF.FinishFunction();
3558 return Fn;
3559}
3560
3561/// This function emits a helper that reduces all the reduction variables from
3562/// the team into the provided global buffer for the reduction variables.
3563///
3564/// void global_to_list_reduce_func(void *buffer, int Idx, void *reduce_data)
3565/// void *GlobPtrs[];
3566/// GlobPtrs[0] = (void*)&buffer.D0[Idx];
3567/// ...
3568/// GlobPtrs[N] = (void*)&buffer.DN[Idx];
3569/// reduce_function(reduce_data, GlobPtrs);
3570static llvm::Value *emitGlobalToListReduceFunction(
3571 CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
3572 QualType ReductionArrayTy, SourceLocation Loc,
3573 const RecordDecl *TeamReductionRec,
3574 const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
3575 &VarFieldMap,
3576 llvm::Function *ReduceFn) {
3577 ASTContext &C = CGM.getContext();
3578
3579 // Buffer: global reduction buffer.
3580 ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3581 C.VoidPtrTy, ImplicitParamDecl::Other);
3582 // Idx: index of the buffer.
3583 ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
3584 ImplicitParamDecl::Other);
3585 // ReduceList: thread local Reduce list.
3586 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
3587 C.VoidPtrTy, ImplicitParamDecl::Other);
3588 FunctionArgList Args;
3589 Args.push_back(&BufferArg);
3590 Args.push_back(&IdxArg);
3591 Args.push_back(&ReduceListArg);
3592
3593 const CGFunctionInfo &CGFI =
3594 CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
3595 auto *Fn = llvm::Function::Create(
3596 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
3597 "_omp_reduction_global_to_list_reduce_func", &CGM.getModule());
3598 CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
3599 Fn->setDoesNotRecurse();
3600 CodeGenFunction CGF(CGM);
3601 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
3602
3603 CGBuilderTy &Bld = CGF.Builder;
3604
3605 Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
3606 QualType StaticTy = C.getRecordType(TeamReductionRec);
3607 llvm::Type *LLVMReductionsBufferTy =
3608 CGM.getTypes().ConvertTypeForMem(StaticTy);
3609 llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
3610 CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
3611 LLVMReductionsBufferTy->getPointerTo());
3612
3613 // 1. Build a list of reduction variables.
3614 // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
3615 Address ReductionList =
3616 CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.red_list");
3617 auto IPriv = Privates.begin();
3618 llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
3619 CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
3620 /*Volatile=*/false, C.IntTy,
3621 Loc)};
3622 unsigned Idx = 0;
3623 for (unsigned I = 0, E = Privates.size(); I < E; ++I, ++IPriv, ++Idx) {
3624 Address Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3625 // Global = Buffer.VD[Idx];
3626 const ValueDecl *VD = cast<DeclRefExpr>(*IPriv)->getDecl();
3627 const FieldDecl *FD = VarFieldMap.lookup(VD);
3628 LValue GlobLVal = CGF.EmitLValueForField(
3629 CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
3630 Address GlobAddr = GlobLVal.getAddress(CGF);
3631 llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
3632 GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
3633 llvm::Value *Ptr = CGF.EmitCastToVoidPtr(BufferPtr);
3634 CGF.EmitStoreOfScalar(Ptr, Elem, /*Volatile=*/false, C.VoidPtrTy);
3635 if ((*IPriv)->getType()->isVariablyModifiedType()) {
3636 // Store array size.
3637 ++Idx;
3638 Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3639 llvm::Value *Size = CGF.Builder.CreateIntCast(
3640 CGF.getVLASize(
3641 CGF.getContext().getAsVariableArrayType((*IPriv)->getType()))
3642 .NumElts,
3643 CGF.SizeTy, /*isSigned=*/false);
3644 CGF.Builder.CreateStore(CGF.Builder.CreateIntToPtr(Size, CGF.VoidPtrTy),
3645 Elem);
3646 }
3647 }
3648
3649 // Call reduce_function(ReduceList, GlobalReduceList)
3650 llvm::Value *GlobalReduceList =
3651 CGF.EmitCastToVoidPtr(ReductionList.getPointer());
3652 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
3653 llvm::Value *ReducedPtr = CGF.EmitLoadOfScalar(
3654 AddrReduceListArg, /*Volatile=*/false, C.VoidPtrTy, Loc);
3655 CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
3656 CGF, Loc, ReduceFn, {ReducedPtr, GlobalReduceList});
3657 CGF.FinishFunction();
3658 return Fn;
3659}
3660
3661///
3662/// Design of OpenMP reductions on the GPU
3663///
3664/// Consider a typical OpenMP program with one or more reduction
3665/// clauses:
3666///
3667/// float foo;
3668/// double bar;
3669/// #pragma omp target teams distribute parallel for \
3670/// reduction(+:foo) reduction(*:bar)
3671/// for (int i = 0; i < N; i++) {
3672/// foo += A[i]; bar *= B[i];
3673/// }
3674///
3675/// where 'foo' and 'bar' are reduced across all OpenMP threads in
3676/// all teams. In our OpenMP implementation on the NVPTX device an
3677/// OpenMP team is mapped to a CUDA threadblock and OpenMP threads
3678/// within a team are mapped to CUDA threads within a threadblock.
3679/// Our goal is to efficiently aggregate values across all OpenMP
3680/// threads such that:
3681///
3682/// - the compiler and runtime are logically concise, and
3683/// - the reduction is performed efficiently in a hierarchical
3684/// manner as follows: within OpenMP threads in the same warp,
3685/// across warps in a threadblock, and finally across teams on
3686/// the NVPTX device.
3687///
3688/// Introduction to Decoupling
3689///
3690/// We would like to decouple the compiler and the runtime so that the
3691/// latter is ignorant of the reduction variables (number, data types)
3692/// and the reduction operators. This allows a simpler interface
3693/// and implementation while still attaining good performance.
3694///
3695/// Pseudocode for the aforementioned OpenMP program generated by the
3696/// compiler is as follows:
3697///
3698/// 1. Create private copies of reduction variables on each OpenMP
3699/// thread: 'foo_private', 'bar_private'
3700/// 2. Each OpenMP thread reduces the chunk of 'A' and 'B' assigned
3701/// to it and writes the result in 'foo_private' and 'bar_private'
3702/// respectively.
3703/// 3. Call the OpenMP runtime on the GPU to reduce within a team
3704/// and store the result on the team master:
3705///
3706/// __kmpc_nvptx_parallel_reduce_nowait_v2(...,
3707/// reduceData, shuffleReduceFn, interWarpCpyFn)
3708///
3709/// where:
3710/// struct ReduceData {
3711/// double *foo;
3712/// double *bar;
3713/// } reduceData
3714/// reduceData.foo = &foo_private
3715/// reduceData.bar = &bar_private
3716///
3717/// 'shuffleReduceFn' and 'interWarpCpyFn' are pointers to two
3718/// auxiliary functions generated by the compiler that operate on
3719/// variables of type 'ReduceData'. They aid the runtime perform
3720/// algorithmic steps in a data agnostic manner.
3721///
3722/// 'shuffleReduceFn' is a pointer to a function that reduces data
3723/// of type 'ReduceData' across two OpenMP threads (lanes) in the
3724/// same warp. It takes the following arguments as input:
3725///
3726/// a. variable of type 'ReduceData' on the calling lane,
3727/// b. its lane_id,
3728/// c. an offset relative to the current lane_id to generate a
3729/// remote_lane_id. The remote lane contains the second
3730/// variable of type 'ReduceData' that is to be reduced.
3731/// d. an algorithm version parameter determining which reduction
3732/// algorithm to use.
3733///
3734/// 'shuffleReduceFn' retrieves data from the remote lane using
3735/// efficient GPU shuffle intrinsics and reduces, using the
3736/// algorithm specified by the 4th parameter, the two operands
3737/// element-wise. The result is written to the first operand.
3738///
3739/// Different reduction algorithms are implemented in different
3740/// runtime functions, all calling 'shuffleReduceFn' to perform
3741/// the essential reduction step. Therefore, based on the 4th
3742/// parameter, this function behaves slightly differently to
3743/// cooperate with the runtime to ensure correctness under
3744/// different circumstances.
3745///
3746/// 'InterWarpCpyFn' is a pointer to a function that transfers
3747/// reduced variables across warps. It tunnels, through CUDA
3748/// shared memory, the thread-private data of type 'ReduceData'
3749/// from lane 0 of each warp to a lane in the first warp.
3750/// 4. Call the OpenMP runtime on the GPU to reduce across teams.
3751/// The last team writes the global reduced value to memory.
3752///
3753/// ret = __kmpc_nvptx_teams_reduce_nowait(...,
3754/// reduceData, shuffleReduceFn, interWarpCpyFn,
3755/// scratchpadCopyFn, loadAndReduceFn)
3756///
3757/// 'scratchpadCopyFn' is a helper that stores reduced
3758/// data from the team master to a scratchpad array in
3759/// global memory.
3760///
3761/// 'loadAndReduceFn' is a helper that loads data from
3762/// the scratchpad array and reduces it with the input
3763/// operand.
3764///
3765/// These compiler generated functions hide address
3766/// calculation and alignment information from the runtime.
3767/// 5. if ret == 1:
3768/// The team master of the last team stores the reduced
3769/// result to the globals in memory.
3770/// foo += reduceData.foo; bar *= reduceData.bar
3771///
3772///
3773/// Warp Reduction Algorithms
3774///
3775/// On the warp level, we have three algorithms implemented in the
3776/// OpenMP runtime depending on the number of active lanes:
3777///
3778/// Full Warp Reduction
3779///
3780/// The reduce algorithm within a warp where all lanes are active
3781/// is implemented in the runtime as follows:
3782///
3783/// full_warp_reduce(void *reduce_data,
3784/// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
3785/// for (int offset = WARPSIZE/2; offset > 0; offset /= 2)
3786/// ShuffleReduceFn(reduce_data, 0, offset, 0);
3787/// }
3788///
3789/// The algorithm completes in log(2, WARPSIZE) steps.
3790///
3791/// 'ShuffleReduceFn' is used here with lane_id set to 0 because it is
3792/// not used therefore we save instructions by not retrieving lane_id
3793/// from the corresponding special registers. The 4th parameter, which
3794/// represents the version of the algorithm being used, is set to 0 to
3795/// signify full warp reduction.
3796///
3797/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
3798///
3799/// #reduce_elem refers to an element in the local lane's data structure
3800/// #remote_elem is retrieved from a remote lane
3801/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
3802/// reduce_elem = reduce_elem REDUCE_OP remote_elem;
3803///
3804/// Contiguous Partial Warp Reduction
3805///
3806/// This reduce algorithm is used within a warp where only the first
3807/// 'n' (n <= WARPSIZE) lanes are active. It is typically used when the
3808/// number of OpenMP threads in a parallel region is not a multiple of
3809/// WARPSIZE. The algorithm is implemented in the runtime as follows:
3810///
3811/// void
3812/// contiguous_partial_reduce(void *reduce_data,
3813/// kmp_ShuffleReductFctPtr ShuffleReduceFn,
3814/// int size, int lane_id) {
3815/// int curr_size;
3816/// int offset;
3817/// curr_size = size;
3818/// mask = curr_size/2;
3819/// while (offset>0) {
3820/// ShuffleReduceFn(reduce_data, lane_id, offset, 1);
3821/// curr_size = (curr_size+1)/2;
3822/// offset = curr_size/2;
3823/// }
3824/// }
3825///
3826/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
3827///
3828/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
3829/// if (lane_id < offset)
3830/// reduce_elem = reduce_elem REDUCE_OP remote_elem
3831/// else
3832/// reduce_elem = remote_elem
3833///
3834/// This algorithm assumes that the data to be reduced are located in a
3835/// contiguous subset of lanes starting from the first. When there is
3836/// an odd number of active lanes, the data in the last lane is not
3837/// aggregated with any other lane's dat but is instead copied over.
3838///
3839/// Dispersed Partial Warp Reduction
3840///
3841/// This algorithm is used within a warp when any discontiguous subset of
3842/// lanes are active. It is used to implement the reduction operation
3843/// across lanes in an OpenMP simd region or in a nested parallel region.
3844///
3845/// void
3846/// dispersed_partial_reduce(void *reduce_data,
3847/// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
3848/// int size, remote_id;
3849/// int logical_lane_id = number_of_active_lanes_before_me() * 2;
3850/// do {
3851/// remote_id = next_active_lane_id_right_after_me();
3852/// # the above function returns 0 of no active lane
3853/// # is present right after the current lane.
3854/// size = number_of_active_lanes_in_this_warp();
3855/// logical_lane_id /= 2;
3856/// ShuffleReduceFn(reduce_data, logical_lane_id,
3857/// remote_id-1-threadIdx.x, 2);
3858/// } while (logical_lane_id % 2 == 0 && size > 1);
3859/// }
3860///
3861/// There is no assumption made about the initial state of the reduction.
3862/// Any number of lanes (>=1) could be active at any position. The reduction
3863/// result is returned in the first active lane.
3864///
3865/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
3866///
3867/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
3868/// if (lane_id % 2 == 0 && offset > 0)
3869/// reduce_elem = reduce_elem REDUCE_OP remote_elem
3870/// else
3871/// reduce_elem = remote_elem
3872///
3873///
3874/// Intra-Team Reduction
3875///
3876/// This function, as implemented in the runtime call
3877/// '__kmpc_nvptx_parallel_reduce_nowait_v2', aggregates data across OpenMP
3878/// threads in a team. It first reduces within a warp using the
3879/// aforementioned algorithms. We then proceed to gather all such
3880/// reduced values at the first warp.
3881///
3882/// The runtime makes use of the function 'InterWarpCpyFn', which copies
3883/// data from each of the "warp master" (zeroth lane of each warp, where
3884/// warp-reduced data is held) to the zeroth warp. This step reduces (in
3885/// a mathematical sense) the problem of reduction across warp masters in
3886/// a block to the problem of warp reduction.
3887///
3888///
3889/// Inter-Team Reduction
3890///
3891/// Once a team has reduced its data to a single value, it is stored in
3892/// a global scratchpad array. Since each team has a distinct slot, this
3893/// can be done without locking.
3894///
3895/// The last team to write to the scratchpad array proceeds to reduce the
3896/// scratchpad array. One or more workers in the last team use the helper
3897/// 'loadAndReduceDataFn' to load and reduce values from the array, i.e.,
3898/// the k'th worker reduces every k'th element.
3899///
3900/// Finally, a call is made to '__kmpc_nvptx_parallel_reduce_nowait_v2' to
3901/// reduce across workers and compute a globally reduced value.
3902///
3903void CGOpenMPRuntimeGPU::emitReduction(
3904 CodeGenFunction &CGF, SourceLocation Loc, ArrayRef<const Expr *> Privates,
3905 ArrayRef<const Expr *> LHSExprs, ArrayRef<const Expr *> RHSExprs,
3906 ArrayRef<const Expr *> ReductionOps, ReductionOptionsTy Options) {
3907 if (!CGF.HaveInsertPoint())
3908 return;
3909
3910 bool ParallelReduction = isOpenMPParallelDirective(Options.ReductionKind);
3911#ifndef NDEBUG
3912 bool TeamsReduction = isOpenMPTeamsDirective(Options.ReductionKind);
3913#endif
3914
3915 if (Options.SimpleReduction) {
3916 assert(!TeamsReduction && !ParallelReduction &&((!TeamsReduction && !ParallelReduction && "Invalid reduction selection in emitReduction."
) ? static_cast<void> (0) : __assert_fail ("!TeamsReduction && !ParallelReduction && \"Invalid reduction selection in emitReduction.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 3917, __PRETTY_FUNCTION__))
3917 "Invalid reduction selection in emitReduction.")((!TeamsReduction && !ParallelReduction && "Invalid reduction selection in emitReduction."
) ? static_cast<void> (0) : __assert_fail ("!TeamsReduction && !ParallelReduction && \"Invalid reduction selection in emitReduction.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 3917, __PRETTY_FUNCTION__))
;
3918 CGOpenMPRuntime::emitReduction(CGF, Loc, Privates, LHSExprs, RHSExprs,
3919 ReductionOps, Options);
3920 return;
3921 }
3922
3923 assert((TeamsReduction || ParallelReduction) &&(((TeamsReduction || ParallelReduction) && "Invalid reduction selection in emitReduction."
) ? static_cast<void> (0) : __assert_fail ("(TeamsReduction || ParallelReduction) && \"Invalid reduction selection in emitReduction.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 3924, __PRETTY_FUNCTION__))
3924 "Invalid reduction selection in emitReduction.")(((TeamsReduction || ParallelReduction) && "Invalid reduction selection in emitReduction."
) ? static_cast<void> (0) : __assert_fail ("(TeamsReduction || ParallelReduction) && \"Invalid reduction selection in emitReduction.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 3924, __PRETTY_FUNCTION__))
;
3925
3926 // Build res = __kmpc_reduce{_nowait}(<gtid>, <n>, sizeof(RedList),
3927 // RedList, shuffle_reduce_func, interwarp_copy_func);
3928 // or
3929 // Build res = __kmpc_reduce_teams_nowait_simple(<loc>, <gtid>, <lck>);
3930 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
3931 llvm::Value *ThreadId = getThreadID(CGF, Loc);
3932
3933 llvm::Value *Res;
3934 ASTContext &C = CGM.getContext();
3935 // 1. Build a list of reduction variables.
3936 // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
3937 auto Size = RHSExprs.size();
3938 for (const Expr *E : Privates) {
3939 if (E->getType()->isVariablyModifiedType())
3940 // Reserve place for array size.
3941 ++Size;
3942 }
3943 llvm::APInt ArraySize(/*unsigned int numBits=*/32, Size);
3944 QualType ReductionArrayTy =
3945 C.getConstantArrayType(C.VoidPtrTy, ArraySize, nullptr, ArrayType::Normal,
3946 /*IndexTypeQuals=*/0);
3947 Address ReductionList =
3948 CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.red_list");
3949 auto IPriv = Privates.begin();
3950 unsigned Idx = 0;
3951 for (unsigned I = 0, E = RHSExprs.size(); I < E; ++I, ++IPriv, ++Idx) {
3952 Address Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3953 CGF.Builder.CreateStore(
3954 CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3955 CGF.EmitLValue(RHSExprs[I]).getPointer(CGF), CGF.VoidPtrTy),
3956 Elem);
3957 if ((*IPriv)->getType()->isVariablyModifiedType()) {
3958 // Store array size.
3959 ++Idx;
3960 Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3961 llvm::Value *Size = CGF.Builder.CreateIntCast(
3962 CGF.getVLASize(
3963 CGF.getContext().getAsVariableArrayType((*IPriv)->getType()))
3964 .NumElts,
3965 CGF.SizeTy, /*isSigned=*/false);
3966 CGF.Builder.CreateStore(CGF.Builder.CreateIntToPtr(Size, CGF.VoidPtrTy),
3967 Elem);
3968 }
3969 }
3970
3971 llvm::Value *RL = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3972 ReductionList.getPointer(), CGF.VoidPtrTy);
3973 llvm::Function *ReductionFn = emitReductionFunction(
3974 Loc, CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo(), Privates,
3975 LHSExprs, RHSExprs, ReductionOps);
3976 llvm::Value *ReductionArrayTySize = CGF.getTypeSize(ReductionArrayTy);
3977 llvm::Function *ShuffleAndReduceFn = emitShuffleAndReduceFunction(
3978 CGM, Privates, ReductionArrayTy, ReductionFn, Loc);
3979 llvm::Value *InterWarpCopyFn =
3980 emitInterWarpCopyFunction(CGM, Privates, ReductionArrayTy, Loc);
3981
3982 if (ParallelReduction) {
3983 llvm::Value *Args[] = {RTLoc,
3984 ThreadId,
3985 CGF.Builder.getInt32(RHSExprs.size()),
3986 ReductionArrayTySize,
3987 RL,
3988 ShuffleAndReduceFn,
3989 InterWarpCopyFn};
3990
3991 Res = CGF.EmitRuntimeCall(
3992 OMPBuilder.getOrCreateRuntimeFunction(
3993 CGM.getModule(), OMPRTL___kmpc_nvptx_parallel_reduce_nowait_v2),
3994 Args);
3995 } else {
3996 assert(TeamsReduction && "expected teams reduction.")((TeamsReduction && "expected teams reduction.") ? static_cast
<void> (0) : __assert_fail ("TeamsReduction && \"expected teams reduction.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 3996, __PRETTY_FUNCTION__))
;
3997 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> VarFieldMap;
3998 llvm::SmallVector<const ValueDecl *, 4> PrivatesReductions(Privates.size());
3999 int Cnt = 0;
4000 for (const Expr *DRE : Privates) {
4001 PrivatesReductions[Cnt] = cast<DeclRefExpr>(DRE)->getDecl();
4002 ++Cnt;
4003 }
4004 const RecordDecl *TeamReductionRec = ::buildRecordForGlobalizedVars(
4005 CGM.getContext(), PrivatesReductions, llvm::None, VarFieldMap,
4006 C.getLangOpts().OpenMPCUDAReductionBufNum);
4007 TeamsReductions.push_back(TeamReductionRec);
4008 if (!KernelTeamsReductionPtr) {
4009 KernelTeamsReductionPtr = new llvm::GlobalVariable(
4010 CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/true,
4011 llvm::GlobalValue::InternalLinkage, nullptr,
4012 "_openmp_teams_reductions_buffer_$_$ptr");
4013 }
4014 llvm::Value *GlobalBufferPtr = CGF.EmitLoadOfScalar(
4015 Address(KernelTeamsReductionPtr, CGM.getPointerAlign()),
4016 /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc);
4017 llvm::Value *GlobalToBufferCpyFn = ::emitListToGlobalCopyFunction(
4018 CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap);
4019 llvm::Value *GlobalToBufferRedFn = ::emitListToGlobalReduceFunction(
4020 CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap,
4021 ReductionFn);
4022 llvm::Value *BufferToGlobalCpyFn = ::emitGlobalToListCopyFunction(
4023 CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap);
4024 llvm::Value *BufferToGlobalRedFn = ::emitGlobalToListReduceFunction(
4025 CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap,
4026 ReductionFn);
4027
4028 llvm::Value *Args[] = {
4029 RTLoc,
4030 ThreadId,
4031 GlobalBufferPtr,
4032 CGF.Builder.getInt32(C.getLangOpts().OpenMPCUDAReductionBufNum),
4033 RL,
4034 ShuffleAndReduceFn,
4035 InterWarpCopyFn,
4036 GlobalToBufferCpyFn,
4037 GlobalToBufferRedFn,
4038 BufferToGlobalCpyFn,
4039 BufferToGlobalRedFn};
4040
4041 Res = CGF.EmitRuntimeCall(
4042 OMPBuilder.getOrCreateRuntimeFunction(
4043 CGM.getModule(), OMPRTL___kmpc_nvptx_teams_reduce_nowait_v2),
4044 Args);
4045 }
4046
4047 // 5. Build if (res == 1)
4048 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".omp.reduction.done");
4049 llvm::BasicBlock *ThenBB = CGF.createBasicBlock(".omp.reduction.then");
4050 llvm::Value *Cond = CGF.Builder.CreateICmpEQ(
4051 Res, llvm::ConstantInt::get(CGM.Int32Ty, /*V=*/1));
4052 CGF.Builder.CreateCondBr(Cond, ThenBB, ExitBB);
4053
4054 // 6. Build then branch: where we have reduced values in the master
4055 // thread in each team.
4056 // __kmpc_end_reduce{_nowait}(<gtid>);
4057 // break;
4058 CGF.EmitBlock(ThenBB);
4059
4060 // Add emission of __kmpc_end_reduce{_nowait}(<gtid>);
4061 auto &&CodeGen = [Privates, LHSExprs, RHSExprs, ReductionOps,
4062 this](CodeGenFunction &CGF, PrePostActionTy &Action) {
4063 auto IPriv = Privates.begin();
4064 auto ILHS = LHSExprs.begin();
4065 auto IRHS = RHSExprs.begin();
4066 for (const Expr *E : ReductionOps) {
4067 emitSingleReductionCombiner(CGF, E, *IPriv, cast<DeclRefExpr>(*ILHS),
4068 cast<DeclRefExpr>(*IRHS));
4069 ++IPriv;
4070 ++ILHS;
4071 ++IRHS;
4072 }
4073 };
4074 llvm::Value *EndArgs[] = {ThreadId};
4075 RegionCodeGenTy RCG(CodeGen);
4076 NVPTXActionTy Action(
4077 nullptr, llvm::None,
4078 OMPBuilder.getOrCreateRuntimeFunction(
4079 CGM.getModule(), OMPRTL___kmpc_nvptx_end_reduce_nowait),
4080 EndArgs);
4081 RCG.setAction(Action);
4082 RCG(CGF);
4083 // There is no need to emit line number for unconditional branch.
4084 (void)ApplyDebugLocation::CreateEmpty(CGF);
4085 CGF.EmitBlock(ExitBB, /*IsFinished=*/true);
4086}
4087
4088const VarDecl *
4089CGOpenMPRuntimeGPU::translateParameter(const FieldDecl *FD,
4090 const VarDecl *NativeParam) const {
4091 if (!NativeParam->getType()->isReferenceType())
4092 return NativeParam;
4093 QualType ArgType = NativeParam->getType();
4094 QualifierCollector QC;
4095 const Type *NonQualTy = QC.strip(ArgType);
4096 QualType PointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
4097 if (const auto *Attr = FD->getAttr<OMPCaptureKindAttr>()) {
4098 if (Attr->getCaptureKind() == OMPC_map) {
4099 PointeeTy = CGM.getContext().getAddrSpaceQualType(PointeeTy,
4100 LangAS::opencl_global);
4101 } else if (Attr->getCaptureKind() == OMPC_firstprivate &&
4102 PointeeTy.isConstant(CGM.getContext())) {
4103 PointeeTy = CGM.getContext().getAddrSpaceQualType(PointeeTy,
4104 LangAS::opencl_generic);
4105 }
4106 }
4107 ArgType = CGM.getContext().getPointerType(PointeeTy);
4108 QC.addRestrict();
4109 enum { NVPTX_local_addr = 5 };
4110 QC.addAddressSpace(getLangASFromTargetAS(NVPTX_local_addr));
4111 ArgType = QC.apply(CGM.getContext(), ArgType);
4112 if (isa<ImplicitParamDecl>(NativeParam))
4113 return ImplicitParamDecl::Create(
4114 CGM.getContext(), /*DC=*/nullptr, NativeParam->getLocation(),
4115 NativeParam->getIdentifier(), ArgType, ImplicitParamDecl::Other);
4116 return ParmVarDecl::Create(
4117 CGM.getContext(),
4118 const_cast<DeclContext *>(NativeParam->getDeclContext()),
4119 NativeParam->getBeginLoc(), NativeParam->getLocation(),
4120 NativeParam->getIdentifier(), ArgType,
4121 /*TInfo=*/nullptr, SC_None, /*DefArg=*/nullptr);
4122}
4123
4124Address
4125CGOpenMPRuntimeGPU::getParameterAddress(CodeGenFunction &CGF,
4126 const VarDecl *NativeParam,
4127 const VarDecl *TargetParam) const {
4128 assert(NativeParam != TargetParam &&((NativeParam != TargetParam && NativeParam->getType
()->isReferenceType() && "Native arg must not be the same as target arg."
) ? static_cast<void> (0) : __assert_fail ("NativeParam != TargetParam && NativeParam->getType()->isReferenceType() && \"Native arg must not be the same as target arg.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4130, __PRETTY_FUNCTION__))
4129 NativeParam->getType()->isReferenceType() &&((NativeParam != TargetParam && NativeParam->getType
()->isReferenceType() && "Native arg must not be the same as target arg."
) ? static_cast<void> (0) : __assert_fail ("NativeParam != TargetParam && NativeParam->getType()->isReferenceType() && \"Native arg must not be the same as target arg.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4130, __PRETTY_FUNCTION__))
4130 "Native arg must not be the same as target arg.")((NativeParam != TargetParam && NativeParam->getType
()->isReferenceType() && "Native arg must not be the same as target arg."
) ? static_cast<void> (0) : __assert_fail ("NativeParam != TargetParam && NativeParam->getType()->isReferenceType() && \"Native arg must not be the same as target arg.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4130, __PRETTY_FUNCTION__))
;
4131 Address LocalAddr = CGF.GetAddrOfLocalVar(TargetParam);
4132 QualType NativeParamType = NativeParam->getType();
4133 QualifierCollector QC;
4134 const Type *NonQualTy = QC.strip(NativeParamType);
4135 QualType NativePointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
4136 unsigned NativePointeeAddrSpace =
4137 CGF.getContext().getTargetAddressSpace(NativePointeeTy);
4138 QualType TargetTy = TargetParam->getType();
4139 llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(
4140 LocalAddr, /*Volatile=*/false, TargetTy, SourceLocation());
4141 // First cast to generic.
4142 TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
4143 TargetAddr, TargetAddr->getType()->getPointerElementType()->getPointerTo(
4144 /*AddrSpace=*/0));
4145 // Cast from generic to native address space.
4146 TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
4147 TargetAddr, TargetAddr->getType()->getPointerElementType()->getPointerTo(
4148 NativePointeeAddrSpace));
4149 Address NativeParamAddr = CGF.CreateMemTemp(NativeParamType);
4150 CGF.EmitStoreOfScalar(TargetAddr, NativeParamAddr, /*Volatile=*/false,
4151 NativeParamType);
4152 return NativeParamAddr;
4153}
4154
4155void CGOpenMPRuntimeGPU::emitOutlinedFunctionCall(
4156 CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn,
4157 ArrayRef<llvm::Value *> Args) const {
4158 SmallVector<llvm::Value *, 4> TargetArgs;
4159 TargetArgs.reserve(Args.size());
4160 auto *FnType = OutlinedFn.getFunctionType();
4161 for (unsigned I = 0, E = Args.size(); I < E; ++I) {
4162 if (FnType->isVarArg() && FnType->getNumParams() <= I) {
4163 TargetArgs.append(std::next(Args.begin(), I), Args.end());
4164 break;
4165 }
4166 llvm::Type *TargetType = FnType->getParamType(I);
4167 llvm::Value *NativeArg = Args[I];
4168 if (!TargetType->isPointerTy()) {
4169 TargetArgs.emplace_back(NativeArg);
4170 continue;
4171 }
4172 llvm::Value *TargetArg = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
4173 NativeArg,
4174 NativeArg->getType()->getPointerElementType()->getPointerTo());
4175 TargetArgs.emplace_back(
4176 CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(TargetArg, TargetType));
4177 }
4178 CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, TargetArgs);
4179}
4180
4181/// Emit function which wraps the outline parallel region
4182/// and controls the arguments which are passed to this function.
4183/// The wrapper ensures that the outlined function is called
4184/// with the correct arguments when data is shared.
4185llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
4186 llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) {
4187 ASTContext &Ctx = CGM.getContext();
4188 const auto &CS = *D.getCapturedStmt(OMPD_parallel);
4189
4190 // Create a function that takes as argument the source thread.
4191 FunctionArgList WrapperArgs;
4192 QualType Int16QTy =
4193 Ctx.getIntTypeForBitwidth(/*DestWidth=*/16, /*Signed=*/false);
4194 QualType Int32QTy =
4195 Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/false);
4196 ImplicitParamDecl ParallelLevelArg(Ctx, /*DC=*/nullptr, D.getBeginLoc(),
4197 /*Id=*/nullptr, Int16QTy,
4198 ImplicitParamDecl::Other);
4199 ImplicitParamDecl WrapperArg(Ctx, /*DC=*/nullptr, D.getBeginLoc(),
4200 /*Id=*/nullptr, Int32QTy,
4201 ImplicitParamDecl::Other);
4202 WrapperArgs.emplace_back(&ParallelLevelArg);
4203 WrapperArgs.emplace_back(&WrapperArg);
4204
4205 const CGFunctionInfo &CGFI =
4206 CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, WrapperArgs);
4207
4208 auto *Fn = llvm::Function::Create(
4209 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
4210 Twine(OutlinedParallelFn->getName(), "_wrapper"), &CGM.getModule());
4211
4212 // Ensure we do not inline the function. This is trivially true for the ones
4213 // passed to __kmpc_fork_call but the ones calles in serialized regions
4214 // could be inlined. This is not a perfect but it is closer to the invariant
4215 // we want, namely, every data environment starts with a new function.
4216 // TODO: We should pass the if condition to the runtime function and do the
4217 // handling there. Much cleaner code.
4218 Fn->addFnAttr(llvm::Attribute::NoInline);
4219
4220 CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
4221 Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
4222 Fn->setDoesNotRecurse();
4223
4224 CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
4225 CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, Fn, CGFI, WrapperArgs,
4226 D.getBeginLoc(), D.getBeginLoc());
4227
4228 const auto *RD = CS.getCapturedRecordDecl();
4229 auto CurField = RD->field_begin();
4230
4231 Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
4232 /*Name=*/".zero.addr");
4233 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
4234 // Get the array of arguments.
4235 SmallVector<llvm::Value *, 8> Args;
4236
4237 Args.emplace_back(CGF.GetAddrOfLocalVar(&WrapperArg).getPointer());
4238 Args.emplace_back(ZeroAddr.getPointer());
4239
4240 CGBuilderTy &Bld = CGF.Builder;
4241 auto CI = CS.capture_begin();
4242
4243 // Use global memory for data sharing.
4244 // Handle passing of global args to workers.
4245 Address GlobalArgs =
4246 CGF.CreateDefaultAlignTempAlloca(CGF.VoidPtrPtrTy, "global_args");
4247 llvm::Value *GlobalArgsPtr = GlobalArgs.getPointer();
4248 llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};
4249 CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
4250 CGM.getModule(), OMPRTL___kmpc_get_shared_variables),
4251 DataSharingArgs);
4252
4253 // Retrieve the shared variables from the list of references returned
4254 // by the runtime. Pass the variables to the outlined function.
4255 Address SharedArgListAddress = Address::invalid();
4256 if (CS.capture_size() > 0 ||
4257 isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {
4258 SharedArgListAddress = CGF.EmitLoadOfPointer(
4259 GlobalArgs, CGF.getContext()
4260 .getPointerType(CGF.getContext().getPointerType(
4261 CGF.getContext().VoidPtrTy))
4262 .castAs<PointerType>());
4263 }
4264 unsigned Idx = 0;
4265 if (isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {
4266 Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);
4267 Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
4268 Src, CGF.SizeTy->getPointerTo());
4269 llvm::Value *LB = CGF.EmitLoadOfScalar(
4270 TypedAddress,
4271 /*Volatile=*/false,
4272 CGF.getContext().getPointerType(CGF.getContext().getSizeType()),
4273 cast<OMPLoopDirective>(D).getLowerBoundVariable()->getExprLoc());
4274 Args.emplace_back(LB);
4275 ++Idx;
4276 Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);
4277 TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
4278 Src, CGF.SizeTy->getPointerTo());
4279 llvm::Value *UB = CGF.EmitLoadOfScalar(
4280 TypedAddress,
4281 /*Volatile=*/false,
4282 CGF.getContext().getPointerType(CGF.getContext().getSizeType()),
4283 cast<OMPLoopDirective>(D).getUpperBoundVariable()->getExprLoc());
4284 Args.emplace_back(UB);
4285 ++Idx;
4286 }
4287 if (CS.capture_size() > 0) {
4288 ASTContext &CGFContext = CGF.getContext();
4289 for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CI, ++CurField) {
4290 QualType ElemTy = CurField->getType();
4291 Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, I + Idx);
4292 Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
4293 Src, CGF.ConvertTypeForMem(CGFContext.getPointerType(ElemTy)));
4294 llvm::Value *Arg = CGF.EmitLoadOfScalar(TypedAddress,
4295 /*Volatile=*/false,
4296 CGFContext.getPointerType(ElemTy),
4297 CI->getLocation());
4298 if (CI->capturesVariableByCopy() &&
4299 !CI->getCapturedVar()->getType()->isAnyPointerType()) {
4300 Arg = castValueToType(CGF, Arg, ElemTy, CGFContext.getUIntPtrType(),
4301 CI->getLocation());
4302 }
4303 Args.emplace_back(Arg);
4304 }
4305 }
4306
4307 emitOutlinedFunctionCall(CGF, D.getBeginLoc(), OutlinedParallelFn, Args);
4308 CGF.FinishFunction();
4309 return Fn;
4310}
4311
4312void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,
4313 const Decl *D) {
4314 if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
4315 return;
4316
4317 assert(D && "Expected function or captured|block decl.")((D && "Expected function or captured|block decl.") ?
static_cast<void> (0) : __assert_fail ("D && \"Expected function or captured|block decl.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4317, __PRETTY_FUNCTION__))
;
4318 assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&((FunctionGlobalizedDecls.count(CGF.CurFn) == 0 && "Function is registered already."
) ? static_cast<void> (0) : __assert_fail ("FunctionGlobalizedDecls.count(CGF.CurFn) == 0 && \"Function is registered already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4319, __PRETTY_FUNCTION__))
4319 "Function is registered already.")((FunctionGlobalizedDecls.count(CGF.CurFn) == 0 && "Function is registered already."
) ? static_cast<void> (0) : __assert_fail ("FunctionGlobalizedDecls.count(CGF.CurFn) == 0 && \"Function is registered already.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4319, __PRETTY_FUNCTION__))
;
4320 assert((!TeamAndReductions.first || TeamAndReductions.first == D) &&(((!TeamAndReductions.first || TeamAndReductions.first == D) &&
"Team is set but not processed.") ? static_cast<void> (
0) : __assert_fail ("(!TeamAndReductions.first || TeamAndReductions.first == D) && \"Team is set but not processed.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4321, __PRETTY_FUNCTION__))
4321 "Team is set but not processed.")(((!TeamAndReductions.first || TeamAndReductions.first == D) &&
"Team is set but not processed.") ? static_cast<void> (
0) : __assert_fail ("(!TeamAndReductions.first || TeamAndReductions.first == D) && \"Team is set but not processed.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4321, __PRETTY_FUNCTION__))
;
4322 const Stmt *Body = nullptr;
4323 bool NeedToDelayGlobalization = false;
4324 if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
4325 Body = FD->getBody();
4326 } else if (const auto *BD = dyn_cast<BlockDecl>(D)) {
4327 Body = BD->getBody();
4328 } else if (const auto *CD = dyn_cast<CapturedDecl>(D)) {
4329 Body = CD->getBody();
4330 NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;
4331 if (NeedToDelayGlobalization &&
4332 getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
4333 return;
4334 }
4335 if (!Body)
4336 return;
4337 CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);
4338 VarChecker.Visit(Body);
4339 const RecordDecl *GlobalizedVarsRecord =
4340 VarChecker.getGlobalizedRecord(IsInTTDRegion);
4341 TeamAndReductions.first = nullptr;
4342 TeamAndReductions.second.clear();
4343 ArrayRef<const ValueDecl *> EscapedVariableLengthDecls =
4344 VarChecker.getEscapedVariableLengthDecls();
4345 if (!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty())
4346 return;
4347 auto I = FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;
4348 I->getSecond().MappedParams =
4349 std::make_unique<CodeGenFunction::OMPMapVars>();
4350 I->getSecond().GlobalRecord = GlobalizedVarsRecord;
4351 I->getSecond().EscapedParameters.insert(
4352 VarChecker.getEscapedParameters().begin(),
4353 VarChecker.getEscapedParameters().end());
4354 I->getSecond().EscapedVariableLengthDecls.append(
4355 EscapedVariableLengthDecls.begin(), EscapedVariableLengthDecls.end());
4356 DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
4357 for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
4358 assert(VD->isCanonicalDecl() && "Expected canonical declaration")((VD->isCanonicalDecl() && "Expected canonical declaration"
) ? static_cast<void> (0) : __assert_fail ("VD->isCanonicalDecl() && \"Expected canonical declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4358, __PRETTY_FUNCTION__))
;
4359 const FieldDecl *FD = VarChecker.getFieldForGlobalizedVar(VD);
4360 Data.insert(std::make_pair(VD, MappedVarData(FD, IsInTTDRegion)));
4361 }
4362 if (!IsInTTDRegion && !NeedToDelayGlobalization && !IsInParallelRegion) {
4363 CheckVarsEscapingDeclContext VarChecker(CGF, llvm::None);
4364 VarChecker.Visit(Body);
4365 I->getSecond().SecondaryGlobalRecord =
4366 VarChecker.getGlobalizedRecord(/*IsInTTDRegion=*/true);
4367 I->getSecond().SecondaryLocalVarData.emplace();
4368 DeclToAddrMapTy &Data = I->getSecond().SecondaryLocalVarData.getValue();
4369 for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
4370 assert(VD->isCanonicalDecl() && "Expected canonical declaration")((VD->isCanonicalDecl() && "Expected canonical declaration"
) ? static_cast<void> (0) : __assert_fail ("VD->isCanonicalDecl() && \"Expected canonical declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4370, __PRETTY_FUNCTION__))
;
4371 const FieldDecl *FD = VarChecker.getFieldForGlobalizedVar(VD);
4372 Data.insert(
4373 std::make_pair(VD, MappedVarData(FD, /*IsInTTDRegion=*/true)));
4374 }
4375 }
4376 if (!NeedToDelayGlobalization) {
4377 emitGenericVarsProlog(CGF, D->getBeginLoc(), /*WithSPMDCheck=*/true);
4378 struct GlobalizationScope final : EHScopeStack::Cleanup {
4379 GlobalizationScope() = default;
4380
4381 void Emit(CodeGenFunction &CGF, Flags flags) override {
4382 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
4383 .emitGenericVarsEpilog(CGF, /*WithSPMDCheck=*/true);
4384 }
4385 };
4386 CGF.EHStack.pushCleanup<GlobalizationScope>(NormalAndEHCleanup);
4387 }
4388}
4389
4390Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
4391 const VarDecl *VD) {
4392 if (VD && VD->hasAttr<OMPAllocateDeclAttr>()) {
1
Assuming 'VD' is null
4393 const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
4394 auto AS = LangAS::Default;
4395 switch (A->getAllocatorType()) {
4396 // Use the default allocator here as by default local vars are
4397 // threadlocal.
4398 case OMPAllocateDeclAttr::OMPNullMemAlloc:
4399 case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
4400 case OMPAllocateDeclAttr::OMPThreadMemAlloc:
4401 case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
4402 case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
4403 // Follow the user decision - use default allocation.
4404 return Address::invalid();
4405 case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
4406 // TODO: implement aupport for user-defined allocators.
4407 return Address::invalid();
4408 case OMPAllocateDeclAttr::OMPConstMemAlloc:
4409 AS = LangAS::cuda_constant;
4410 break;
4411 case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
4412 AS = LangAS::cuda_shared;
4413 break;
4414 case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
4415 case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
4416 break;
4417 }
4418 llvm::Type *VarTy = CGF.ConvertTypeForMem(VD->getType());
4419 auto *GV = new llvm::GlobalVariable(
4420 CGM.getModule(), VarTy, /*isConstant=*/false,
4421 llvm::GlobalValue::InternalLinkage, llvm::Constant::getNullValue(VarTy),
4422 VD->getName(),
4423 /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
4424 CGM.getContext().getTargetAddressSpace(AS));
4425 CharUnits Align = CGM.getContext().getDeclAlign(VD);
4426 GV->setAlignment(Align.getAsAlign());
4427 return Address(
4428 CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
4429 GV, VarTy->getPointerTo(CGM.getContext().getTargetAddressSpace(
4430 VD->getType().getAddressSpace()))),
4431 Align);
4432 }
4433
4434 if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
2
Taking false branch
4435 return Address::invalid();
4436
4437 VD = VD->getCanonicalDecl();
3
Called C++ object pointer is null
4438 auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
4439 if (I == FunctionGlobalizedDecls.end())
4440 return Address::invalid();
4441 auto VDI = I->getSecond().LocalVarData.find(VD);
4442 if (VDI != I->getSecond().LocalVarData.end())
4443 return VDI->second.PrivateAddr;
4444 if (VD->hasAttrs()) {
4445 for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),
4446 E(VD->attr_end());
4447 IT != E; ++IT) {
4448 auto VDI = I->getSecond().LocalVarData.find(
4449 cast<VarDecl>(cast<DeclRefExpr>(IT->getRef())->getDecl())
4450 ->getCanonicalDecl());
4451 if (VDI != I->getSecond().LocalVarData.end())
4452 return VDI->second.PrivateAddr;
4453 }
4454 }
4455
4456 return Address::invalid();
4457}
4458
4459void CGOpenMPRuntimeGPU::functionFinished(CodeGenFunction &CGF) {
4460 FunctionGlobalizedDecls.erase(CGF.CurFn);
4461 CGOpenMPRuntime::functionFinished(CGF);
4462}
4463
4464void CGOpenMPRuntimeGPU::getDefaultDistScheduleAndChunk(
4465 CodeGenFunction &CGF, const OMPLoopDirective &S,
4466 OpenMPDistScheduleClauseKind &ScheduleKind,
4467 llvm::Value *&Chunk) const {
4468 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
4469 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
4470 ScheduleKind = OMPC_DIST_SCHEDULE_static;
4471 Chunk = CGF.EmitScalarConversion(
4472 RT.getGPUNumThreads(CGF),
4473 CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
4474 S.getIterationVariable()->getType(), S.getBeginLoc());
4475 return;
4476 }
4477 CGOpenMPRuntime::getDefaultDistScheduleAndChunk(
4478 CGF, S, ScheduleKind, Chunk);
4479}
4480
4481void CGOpenMPRuntimeGPU::getDefaultScheduleAndChunk(
4482 CodeGenFunction &CGF, const OMPLoopDirective &S,
4483 OpenMPScheduleClauseKind &ScheduleKind,
4484 const Expr *&ChunkExpr) const {
4485 ScheduleKind = OMPC_SCHEDULE_static;
4486 // Chunk size is 1 in this case.
4487 llvm::APInt ChunkSize(32, 1);
4488 ChunkExpr = IntegerLiteral::Create(CGF.getContext(), ChunkSize,
4489 CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
4490 SourceLocation());
4491}
4492
4493void CGOpenMPRuntimeGPU::adjustTargetSpecificDataForLambdas(
4494 CodeGenFunction &CGF, const OMPExecutableDirective &D) const {
4495 assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&((isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
" Expected target-based directive.") ? static_cast<void>
(0) : __assert_fail ("isOpenMPTargetExecutionDirective(D.getDirectiveKind()) && \" Expected target-based directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4496, __PRETTY_FUNCTION__))
4496 " Expected target-based directive.")((isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
" Expected target-based directive.") ? static_cast<void>
(0) : __assert_fail ("isOpenMPTargetExecutionDirective(D.getDirectiveKind()) && \" Expected target-based directive.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4496, __PRETTY_FUNCTION__))
;
4497 const CapturedStmt *CS = D.getCapturedStmt(OMPD_target);
4498 for (const CapturedStmt::Capture &C : CS->captures()) {
4499 // Capture variables captured by reference in lambdas for target-based
4500 // directives.
4501 if (!C.capturesVariable())
4502 continue;
4503 const VarDecl *VD = C.getCapturedVar();
4504 const auto *RD = VD->getType()
4505 .getCanonicalType()
4506 .getNonReferenceType()
4507 ->getAsCXXRecordDecl();
4508 if (!RD || !RD->isLambda())
4509 continue;
4510 Address VDAddr = CGF.GetAddrOfLocalVar(VD);
4511 LValue VDLVal;
4512 if (VD->getType().getCanonicalType()->isReferenceType())
4513 VDLVal = CGF.EmitLoadOfReferenceLValue(VDAddr, VD->getType());
4514 else
4515 VDLVal = CGF.MakeAddrLValue(
4516 VDAddr, VD->getType().getCanonicalType().getNonReferenceType());
4517 llvm::DenseMap<const VarDecl *, FieldDecl *> Captures;
4518 FieldDecl *ThisCapture = nullptr;
4519 RD->getCaptureFields(Captures, ThisCapture);
4520 if (ThisCapture && CGF.CapturedStmtInfo->isCXXThisExprCaptured()) {
4521 LValue ThisLVal =
4522 CGF.EmitLValueForFieldInitialization(VDLVal, ThisCapture);
4523 llvm::Value *CXXThis = CGF.LoadCXXThis();
4524 CGF.EmitStoreOfScalar(CXXThis, ThisLVal);
4525 }
4526 for (const LambdaCapture &LC : RD->captures()) {
4527 if (LC.getCaptureKind() != LCK_ByRef)
4528 continue;
4529 const VarDecl *VD = LC.getCapturedVar();
4530 if (!CS->capturesVariable(VD))
4531 continue;
4532 auto It = Captures.find(VD);
4533 assert(It != Captures.end() && "Found lambda capture without field.")((It != Captures.end() && "Found lambda capture without field."
) ? static_cast<void> (0) : __assert_fail ("It != Captures.end() && \"Found lambda capture without field.\""
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4533, __PRETTY_FUNCTION__))
;
4534 LValue VarLVal = CGF.EmitLValueForFieldInitialization(VDLVal, It->second);
4535 Address VDAddr = CGF.GetAddrOfLocalVar(VD);
4536 if (VD->getType().getCanonicalType()->isReferenceType())
4537 VDAddr = CGF.EmitLoadOfReferenceLValue(VDAddr,
4538 VD->getType().getCanonicalType())
4539 .getAddress(CGF);
4540 CGF.EmitStoreOfScalar(VDAddr.getPointer(), VarLVal);
4541 }
4542 }
4543}
4544
4545unsigned CGOpenMPRuntimeGPU::getDefaultFirstprivateAddressSpace() const {
4546 return CGM.getContext().getTargetAddressSpace(LangAS::cuda_constant);
4547}
4548
4549bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
4550 LangAS &AS) {
4551 if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
4552 return false;
4553 const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
4554 switch(A->getAllocatorType()) {
4555 case OMPAllocateDeclAttr::OMPNullMemAlloc:
4556 case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
4557 // Not supported, fallback to the default mem space.
4558 case OMPAllocateDeclAttr::OMPThreadMemAlloc:
4559 case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
4560 case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
4561 case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
4562 case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
4563 AS = LangAS::Default;
4564 return true;
4565 case OMPAllocateDeclAttr::OMPConstMemAlloc:
4566 AS = LangAS::cuda_constant;
4567 return true;
4568 case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
4569 AS = LangAS::cuda_shared;
4570 return true;
4571 case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
4572 llvm_unreachable("Expected predefined allocator for the variables with the "::llvm::llvm_unreachable_internal("Expected predefined allocator for the variables with the "
"static storage.", "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4573)
4573 "static storage.")::llvm::llvm_unreachable_internal("Expected predefined allocator for the variables with the "
"static storage.", "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4573)
;
4574 }
4575 return false;
4576}
4577
4578// Get current CudaArch and ignore any unknown values
4579static CudaArch getCudaArch(CodeGenModule &CGM) {
4580 if (!CGM.getTarget().hasFeature("ptx"))
4581 return CudaArch::UNKNOWN;
4582 for (const auto &Feature : CGM.getTarget().getTargetOpts().FeatureMap) {
4583 if (Feature.getValue()) {
4584 CudaArch Arch = StringToCudaArch(Feature.getKey());
4585 if (Arch != CudaArch::UNKNOWN)
4586 return Arch;
4587 }
4588 }
4589 return CudaArch::UNKNOWN;
4590}
4591
4592/// Check to see if target architecture supports unified addressing which is
4593/// a restriction for OpenMP requires clause "unified_shared_memory".
4594void CGOpenMPRuntimeGPU::processRequiresDirective(
4595 const OMPRequiresDecl *D) {
4596 for (const OMPClause *Clause : D->clauselists()) {
4597 if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
4598 CudaArch Arch = getCudaArch(CGM);
4599 switch (Arch) {
4600 case CudaArch::SM_20:
4601 case CudaArch::SM_21:
4602 case CudaArch::SM_30:
4603 case CudaArch::SM_32:
4604 case CudaArch::SM_35:
4605 case CudaArch::SM_37:
4606 case CudaArch::SM_50:
4607 case CudaArch::SM_52:
4608 case CudaArch::SM_53:
4609 case CudaArch::SM_60:
4610 case CudaArch::SM_61:
4611 case CudaArch::SM_62: {
4612 SmallString<256> Buffer;
4613 llvm::raw_svector_ostream Out(Buffer);
4614 Out << "Target architecture " << CudaArchToString(Arch)
4615 << " does not support unified addressing";
4616 CGM.Error(Clause->getBeginLoc(), Out.str());
4617 return;
4618 }
4619 case CudaArch::SM_70:
4620 case CudaArch::SM_72:
4621 case CudaArch::SM_75:
4622 case CudaArch::SM_80:
4623 case CudaArch::SM_86:
4624 case CudaArch::GFX600:
4625 case CudaArch::GFX601:
4626 case CudaArch::GFX602:
4627 case CudaArch::GFX700:
4628 case CudaArch::GFX701:
4629 case CudaArch::GFX702:
4630 case CudaArch::GFX703:
4631 case CudaArch::GFX704:
4632 case CudaArch::GFX705:
4633 case CudaArch::GFX801:
4634 case CudaArch::GFX802:
4635 case CudaArch::GFX803:
4636 case CudaArch::GFX805:
4637 case CudaArch::GFX810:
4638 case CudaArch::GFX900:
4639 case CudaArch::GFX902:
4640 case CudaArch::GFX904:
4641 case CudaArch::GFX906:
4642 case CudaArch::GFX908:
4643 case CudaArch::GFX909:
4644 case CudaArch::GFX90a:
4645 case CudaArch::GFX90c:
4646 case CudaArch::GFX1010:
4647 case CudaArch::GFX1011:
4648 case CudaArch::GFX1012:
4649 case CudaArch::GFX1030:
4650 case CudaArch::GFX1031:
4651 case CudaArch::GFX1032:
4652 case CudaArch::GFX1033:
4653 case CudaArch::UNUSED:
4654 case CudaArch::UNKNOWN:
4655 break;
4656 case CudaArch::LAST:
4657 llvm_unreachable("Unexpected Cuda arch.")::llvm::llvm_unreachable_internal("Unexpected Cuda arch.", "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4657)
;
4658 }
4659 }
4660 }
4661 CGOpenMPRuntime::processRequiresDirective(D);
4662}
4663
4664/// Get number of SMs and number of blocks per SM.
4665static std::pair<unsigned, unsigned> getSMsBlocksPerSM(CodeGenModule &CGM) {
4666 std::pair<unsigned, unsigned> Data;
4667 if (CGM.getLangOpts().OpenMPCUDANumSMs)
4668 Data.first = CGM.getLangOpts().OpenMPCUDANumSMs;
4669 if (CGM.getLangOpts().OpenMPCUDABlocksPerSM)
4670 Data.second = CGM.getLangOpts().OpenMPCUDABlocksPerSM;
4671 if (Data.first && Data.second)
4672 return Data;
4673 switch (getCudaArch(CGM)) {
4674 case CudaArch::SM_20:
4675 case CudaArch::SM_21:
4676 case CudaArch::SM_30:
4677 case CudaArch::SM_32:
4678 case CudaArch::SM_35:
4679 case CudaArch::SM_37:
4680 case CudaArch::SM_50:
4681 case CudaArch::SM_52:
4682 case CudaArch::SM_53:
4683 return {16, 16};
4684 case CudaArch::SM_60:
4685 case CudaArch::SM_61:
4686 case CudaArch::SM_62:
4687 return {56, 32};
4688 case CudaArch::SM_70:
4689 case CudaArch::SM_72:
4690 case CudaArch::SM_75:
4691 case CudaArch::SM_80:
4692 case CudaArch::SM_86:
4693 return {84, 32};
4694 case CudaArch::GFX600:
4695 case CudaArch::GFX601:
4696 case CudaArch::GFX602:
4697 case CudaArch::GFX700:
4698 case CudaArch::GFX701:
4699 case CudaArch::GFX702:
4700 case CudaArch::GFX703:
4701 case CudaArch::GFX704:
4702 case CudaArch::GFX705:
4703 case CudaArch::GFX801:
4704 case CudaArch::GFX802:
4705 case CudaArch::GFX803:
4706 case CudaArch::GFX805:
4707 case CudaArch::GFX810:
4708 case CudaArch::GFX900:
4709 case CudaArch::GFX902:
4710 case CudaArch::GFX904:
4711 case CudaArch::GFX906:
4712 case CudaArch::GFX908:
4713 case CudaArch::GFX909:
4714 case CudaArch::GFX90a:
4715 case CudaArch::GFX90c:
4716 case CudaArch::GFX1010:
4717 case CudaArch::GFX1011:
4718 case CudaArch::GFX1012:
4719 case CudaArch::GFX1030:
4720 case CudaArch::GFX1031:
4721 case CudaArch::GFX1032:
4722 case CudaArch::GFX1033:
4723 case CudaArch::UNUSED:
4724 case CudaArch::UNKNOWN:
4725 break;
4726 case CudaArch::LAST:
4727 llvm_unreachable("Unexpected Cuda arch.")::llvm::llvm_unreachable_internal("Unexpected Cuda arch.", "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4727)
;
4728 }
4729 llvm_unreachable("Unexpected NVPTX target without ptx feature.")::llvm::llvm_unreachable_internal("Unexpected NVPTX target without ptx feature."
, "/build/llvm-toolchain-snapshot-13~++20210413100635+64c24f493e5f/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp"
, 4729)
;
4730}
4731
4732void CGOpenMPRuntimeGPU::clear() {
4733 if (!GlobalizedRecords.empty() &&
4734 !CGM.getLangOpts().OpenMPCUDATargetParallel) {
4735 ASTContext &C = CGM.getContext();
4736 llvm::SmallVector<const GlobalPtrSizeRecsTy *, 4> GlobalRecs;
4737 llvm::SmallVector<const GlobalPtrSizeRecsTy *, 4> SharedRecs;
4738 RecordDecl *StaticRD = C.buildImplicitRecord(
4739 "_openmp_static_memory_type_$_", RecordDecl::TagKind::TTK_Union);
4740 StaticRD->startDefinition();
4741 RecordDecl *SharedStaticRD = C.buildImplicitRecord(
4742 "_shared_openmp_static_memory_type_$_", RecordDecl::TagKind::TTK_Union);
4743 SharedStaticRD->startDefinition();
4744 for (const GlobalPtrSizeRecsTy &Records : GlobalizedRecords) {
4745 if (Records.Records.empty())
4746 continue;
4747 unsigned Size = 0;
4748 unsigned RecAlignment = 0;
4749 for (const RecordDecl *RD : Records.Records) {
4750 QualType RDTy = C.getRecordType(RD);
4751 unsigned Alignment = C.getTypeAlignInChars(RDTy).getQuantity();
4752 RecAlignment = std::max(RecAlignment, Alignment);
4753 unsigned RecSize = C.getTypeSizeInChars(RDTy).getQuantity();
4754 Size =
4755 llvm::alignTo(llvm::alignTo(Size, Alignment) + RecSize, Alignment);
4756 }
4757 Size = llvm::alignTo(Size, RecAlignment);
4758 llvm::APInt ArySize(/*numBits=*/64, Size);
4759 QualType SubTy = C.getConstantArrayType(
4760 C.CharTy, ArySize, nullptr, ArrayType::Normal, /*IndexTypeQuals=*/0);
4761 const bool UseSharedMemory = Size <= SharedMemorySize;
4762 auto *Field =
4763 FieldDecl::Create(C, UseSharedMemory ? SharedStaticRD : StaticRD,
4764 SourceLocation(), SourceLocation(), nullptr, SubTy,
4765 C.getTrivialTypeSourceInfo(SubTy, SourceLocation()),
4766 /*BW=*/nullptr, /*Mutable=*/false,
4767 /*InitStyle=*/ICIS_NoInit);
4768 Field->setAccess(AS_public);
4769 if (UseSharedMemory) {
4770 SharedStaticRD->addDecl(Field);
4771 SharedRecs.push_back(&Records);
4772 } else {
4773 StaticRD->addDecl(Field);
4774 GlobalRecs.push_back(&Records);
4775 }
4776 Records.RecSize->setInitializer(llvm::ConstantInt::get(CGM.SizeTy, Size));
4777 Records.UseSharedMemory->setInitializer(
4778 llvm::ConstantInt::get(CGM.Int16Ty, UseSharedMemory ? 1 : 0));
4779 }
4780 // Allocate SharedMemorySize buffer for the shared memory.
4781 // FIXME: nvlink does not handle weak linkage correctly (object with the
4782 // different size are reported as erroneous).
4783 // Restore this code as sson as nvlink is fixed.
4784 if (!SharedStaticRD->field_empty()) {
4785 llvm::APInt ArySize(/*numBits=*/64, SharedMemorySize);
4786 QualType SubTy = C.getConstantArrayType(
4787 C.CharTy, ArySize, nullptr, ArrayType::Normal, /*IndexTypeQuals=*/0);
4788 auto *Field = FieldDecl::Create(
4789 C, SharedStaticRD, SourceLocation(), SourceLocation(), nullptr, SubTy,
4790 C.getTrivialTypeSourceInfo(SubTy, SourceLocation()),
4791 /*BW=*/nullptr, /*Mutable=*/false,
4792 /*InitStyle=*/ICIS_NoInit);
4793 Field->setAccess(AS_public);
4794 SharedStaticRD->addDecl(Field);
4795 }
4796 SharedStaticRD->completeDefinition();
4797 if (!SharedStaticRD->field_empty()) {
4798 QualType StaticTy = C.getRecordType(SharedStaticRD);
4799 llvm::Type *LLVMStaticTy = CGM.getTypes().ConvertTypeForMem(StaticTy);
4800 auto *GV = new llvm::GlobalVariable(
4801 CGM.getModule(), LLVMStaticTy,
4802 /*isConstant=*/false, llvm::GlobalValue::WeakAnyLinkage,
4803 llvm::UndefValue::get(LLVMStaticTy),
4804 "_openmp_shared_static_glob_rd_$_", /*InsertBefore=*/nullptr,
4805 llvm::GlobalValue::NotThreadLocal,
4806 C.getTargetAddressSpace(LangAS::cuda_shared));
4807 auto *Replacement = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
4808 GV, CGM.VoidPtrTy);
4809 for (const GlobalPtrSizeRecsTy *Rec : SharedRecs) {
4810 Rec->Buffer->replaceAllUsesWith(Replacement);
4811 Rec->Buffer->eraseFromParent();
4812 }
4813 }
4814 StaticRD->completeDefinition();
4815 if (!StaticRD->field_empty()) {
4816 QualType StaticTy = C.getRecordType(StaticRD);
4817 std::pair<unsigned, unsigned> SMsBlockPerSM = getSMsBlocksPerSM(CGM);
4818 llvm::APInt Size1(32, SMsBlockPerSM.second);
4819 QualType Arr1Ty =
4820 C.getConstantArrayType(StaticTy, Size1, nullptr, ArrayType::Normal,
4821 /*IndexTypeQuals=*/0);
4822 llvm::APInt Size2(32, SMsBlockPerSM.first);
4823 QualType Arr2Ty =
4824 C.getConstantArrayType(Arr1Ty, Size2, nullptr, ArrayType::Normal,
4825 /*IndexTypeQuals=*/0);
4826 llvm::Type *LLVMArr2Ty = CGM.getTypes().ConvertTypeForMem(Arr2Ty);
4827 // FIXME: nvlink does not handle weak linkage correctly (object with the
4828 // different size are reported as erroneous).
4829 // Restore CommonLinkage as soon as nvlink is fixed.
4830 auto *GV = new llvm::GlobalVariable(
4831 CGM.getModule(), LLVMArr2Ty,
4832 /*isConstant=*/false, llvm::GlobalValue::InternalLinkage,
4833 llvm::Constant::getNullValue(LLVMArr2Ty),
4834 "_openmp_static_glob_rd_$_");
4835 auto *Replacement = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
4836 GV, CGM.VoidPtrTy);
4837 for (const GlobalPtrSizeRecsTy *Rec : GlobalRecs) {
4838 Rec->Buffer->replaceAllUsesWith(Replacement);
4839 Rec->Buffer->eraseFromParent();
4840 }
4841 }
4842 }
4843 if (!TeamsReductions.empty()) {
4844 ASTContext &C = CGM.getContext();
4845 RecordDecl *StaticRD = C.buildImplicitRecord(
4846 "_openmp_teams_reduction_type_$_", RecordDecl::TagKind::TTK_Union);
4847 StaticRD->startDefinition();
4848 for (const RecordDecl *TeamReductionRec : TeamsReductions) {
4849 QualType RecTy = C.getRecordType(TeamReductionRec);
4850 auto *Field = FieldDecl::Create(
4851 C, StaticRD, SourceLocation(), SourceLocation(), nullptr, RecTy,
4852 C.getTrivialTypeSourceInfo(RecTy, SourceLocation()),
4853 /*BW=*/nullptr, /*Mutable=*/false,
4854 /*InitStyle=*/ICIS_NoInit);
4855 Field->setAccess(AS_public);
4856 StaticRD->addDecl(Field);
4857 }
4858 StaticRD->completeDefinition();
4859 QualType StaticTy = C.getRecordType(StaticRD);
4860 llvm::Type *LLVMReductionsBufferTy =
4861 CGM.getTypes().ConvertTypeForMem(StaticTy);
4862 // FIXME: nvlink does not handle weak linkage correctly (object with the
4863 // different size are reported as erroneous).
4864 // Restore CommonLinkage as soon as nvlink is fixed.
4865 auto *GV = new llvm::GlobalVariable(
4866 CGM.getModule(), LLVMReductionsBufferTy,
4867 /*isConstant=*/false, llvm::GlobalValue::InternalLinkage,
4868 llvm::Constant::getNullValue(LLVMReductionsBufferTy),
4869 "_openmp_teams_reductions_buffer_$_");
4870 KernelTeamsReductionPtr->setInitializer(
4871 llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(GV,
4872 CGM.VoidPtrTy));
4873 }
4874 CGOpenMPRuntime::clear();
4875}