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, AMDGCN and SPIR-V.
11//
12//===----------------------------------------------------------------------===//
13
14#include "CGOpenMPRuntimeGPU.h"
15#include "CGDebugInfo.h"
16#include "CodeGenFunction.h"
17#include "TargetInfo.h"
18#include "clang/AST/Attr.h"
19#include "clang/AST/DeclOpenMP.h"
20#include "clang/AST/OpenMPClause.h"
21#include "clang/AST/StmtOpenMP.h"
22#include "clang/AST/StmtVisitor.h"
23#include "llvm/ADT/SmallPtrSet.h"
24#include "llvm/Frontend/OpenMP/OMPDeviceConstants.h"
25#include "llvm/Frontend/OpenMP/OMPGridValues.h"
26#include "llvm/IR/IRBuilder.h"
27#include "llvm/IR/Instructions.h"
28#include "llvm/TargetParser/NVPTXTargetParser.h"
29
30using namespace clang;
31using namespace CodeGen;
32using namespace llvm::omp;
33
34namespace {
35/// Pre(post)-action for different OpenMP constructs specialized for NVPTX.
36class NVPTXActionTy final : public PrePostActionTy {
37 llvm::FunctionCallee EnterCallee = nullptr;
38 ArrayRef<llvm::Value *> EnterArgs;
39 llvm::FunctionCallee ExitCallee = nullptr;
40 ArrayRef<llvm::Value *> ExitArgs;
41 bool Conditional = false;
42 llvm::BasicBlock *ContBlock = nullptr;
43
44public:
45 NVPTXActionTy(llvm::FunctionCallee EnterCallee,
46 ArrayRef<llvm::Value *> EnterArgs,
47 llvm::FunctionCallee ExitCallee,
48 ArrayRef<llvm::Value *> ExitArgs, bool Conditional = false)
49 : EnterCallee(EnterCallee), EnterArgs(EnterArgs), ExitCallee(ExitCallee),
50 ExitArgs(ExitArgs), Conditional(Conditional) {}
51 void Enter(CodeGenFunction &CGF) override {
52 llvm::Value *EnterRes = CGF.EmitRuntimeCall(callee: EnterCallee, args: EnterArgs);
53 if (Conditional) {
54 llvm::Value *CallBool = CGF.Builder.CreateIsNotNull(Arg: EnterRes);
55 auto *ThenBlock = CGF.createBasicBlock(name: "omp_if.then");
56 ContBlock = CGF.createBasicBlock(name: "omp_if.end");
57 // Generate the branch (If-stmt)
58 CGF.Builder.CreateCondBr(Cond: CallBool, True: ThenBlock, False: ContBlock);
59 CGF.EmitBlock(BB: ThenBlock);
60 }
61 }
62 void Done(CodeGenFunction &CGF) {
63 // Emit the rest of blocks/branches
64 CGF.EmitBranch(Block: ContBlock);
65 CGF.EmitBlock(BB: ContBlock, IsFinished: true);
66 }
67 void Exit(CodeGenFunction &CGF) override {
68 CGF.EmitRuntimeCall(callee: ExitCallee, args: ExitArgs);
69 }
70};
71
72/// A class to track the execution mode when codegening directives within
73/// a target region. The appropriate mode (SPMD|NON-SPMD) is set on entry
74/// to the target region and used by containing directives such as 'parallel'
75/// to emit optimized code.
76class ExecutionRuntimeModesRAII {
77private:
78 CGOpenMPRuntimeGPU::ExecutionMode SavedExecMode =
79 CGOpenMPRuntimeGPU::EM_Unknown;
80 CGOpenMPRuntimeGPU::ExecutionMode &ExecMode;
81
82public:
83 ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode,
84 CGOpenMPRuntimeGPU::ExecutionMode EntryMode)
85 : ExecMode(ExecMode) {
86 SavedExecMode = ExecMode;
87 ExecMode = EntryMode;
88 }
89 ~ExecutionRuntimeModesRAII() { ExecMode = SavedExecMode; }
90};
91
92static const ValueDecl *getPrivateItem(const Expr *RefExpr) {
93 RefExpr = RefExpr->IgnoreParens();
94 if (const auto *ASE = dyn_cast<ArraySubscriptExpr>(Val: RefExpr)) {
95 const Expr *Base = ASE->getBase()->IgnoreParenImpCasts();
96 while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Val: Base))
97 Base = TempASE->getBase()->IgnoreParenImpCasts();
98 RefExpr = Base;
99 } else if (auto *OASE = dyn_cast<ArraySectionExpr>(Val: RefExpr)) {
100 const Expr *Base = OASE->getBase()->IgnoreParenImpCasts();
101 while (const auto *TempOASE = dyn_cast<ArraySectionExpr>(Val: Base))
102 Base = TempOASE->getBase()->IgnoreParenImpCasts();
103 while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Val: Base))
104 Base = TempASE->getBase()->IgnoreParenImpCasts();
105 RefExpr = Base;
106 }
107 RefExpr = RefExpr->IgnoreParenImpCasts();
108 if (const auto *DE = dyn_cast<DeclRefExpr>(Val: RefExpr))
109 return cast<ValueDecl>(Val: DE->getDecl()->getCanonicalDecl());
110 const auto *ME = cast<MemberExpr>(Val: RefExpr);
111 return cast<ValueDecl>(Val: ME->getMemberDecl()->getCanonicalDecl());
112}
113
114static RecordDecl *buildRecordForGlobalizedVars(
115 ASTContext &C, ArrayRef<const ValueDecl *> EscapedDecls,
116 ArrayRef<const ValueDecl *> EscapedDeclsForTeams,
117 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
118 &MappedDeclsFields,
119 int BufSize) {
120 using VarsDataTy = std::pair<CharUnits /*Align*/, const ValueDecl *>;
121 if (EscapedDecls.empty() && EscapedDeclsForTeams.empty())
122 return nullptr;
123 SmallVector<VarsDataTy, 4> GlobalizedVars;
124 for (const ValueDecl *D : EscapedDecls)
125 GlobalizedVars.emplace_back(Args: C.getDeclAlign(D), Args&: D);
126 for (const ValueDecl *D : EscapedDeclsForTeams)
127 GlobalizedVars.emplace_back(Args: C.getDeclAlign(D), Args&: D);
128
129 // Build struct _globalized_locals_ty {
130 // /* globalized vars */[WarSize] align (decl_align)
131 // /* globalized vars */ for EscapedDeclsForTeams
132 // };
133 RecordDecl *GlobalizedRD = C.buildImplicitRecord(Name: "_globalized_locals_ty");
134 GlobalizedRD->startDefinition();
135 llvm::SmallPtrSet<const ValueDecl *, 16> SingleEscaped(llvm::from_range,
136 EscapedDeclsForTeams);
137 for (const auto &Pair : GlobalizedVars) {
138 const ValueDecl *VD = Pair.second;
139 QualType Type = VD->getType();
140 if (Type->isLValueReferenceType())
141 Type = C.getPointerType(T: Type.getNonReferenceType());
142 else
143 Type = Type.getNonReferenceType();
144 SourceLocation Loc = VD->getLocation();
145 FieldDecl *Field;
146 if (SingleEscaped.count(Ptr: VD)) {
147 Field = FieldDecl::Create(
148 C, DC: GlobalizedRD, StartLoc: Loc, IdLoc: Loc, Id: VD->getIdentifier(), T: Type,
149 TInfo: C.getTrivialTypeSourceInfo(T: Type, Loc: SourceLocation()),
150 /*BW=*/nullptr, /*Mutable=*/false,
151 /*InitStyle=*/ICIS_NoInit);
152 Field->setAccess(AS_public);
153 if (VD->hasAttrs()) {
154 for (specific_attr_iterator<AlignedAttr> I(VD->getAttrs().begin()),
155 E(VD->getAttrs().end());
156 I != E; ++I)
157 Field->addAttr(A: *I);
158 }
159 } else {
160 if (BufSize > 1) {
161 llvm::APInt ArraySize(32, BufSize);
162 Type = C.getConstantArrayType(EltTy: Type, ArySize: ArraySize, SizeExpr: nullptr,
163 ASM: ArraySizeModifier::Normal, IndexTypeQuals: 0);
164 }
165 Field = FieldDecl::Create(
166 C, DC: GlobalizedRD, StartLoc: Loc, IdLoc: Loc, Id: VD->getIdentifier(), T: Type,
167 TInfo: C.getTrivialTypeSourceInfo(T: Type, Loc: SourceLocation()),
168 /*BW=*/nullptr, /*Mutable=*/false,
169 /*InitStyle=*/ICIS_NoInit);
170 Field->setAccess(AS_public);
171 llvm::APInt Align(32, Pair.first.getQuantity());
172 Field->addAttr(A: AlignedAttr::CreateImplicit(
173 Ctx&: C, /*IsAlignmentExpr=*/true,
174 Alignment: IntegerLiteral::Create(C, V: Align,
175 type: C.getIntTypeForBitwidth(DestWidth: 32, /*Signed=*/0),
176 l: SourceLocation()),
177 Range: {}, S: AlignedAttr::GNU_aligned));
178 }
179 GlobalizedRD->addDecl(D: Field);
180 MappedDeclsFields.try_emplace(Key: VD, Args&: Field);
181 }
182 GlobalizedRD->completeDefinition();
183 return GlobalizedRD;
184}
185
186/// Get the list of variables that can escape their declaration context.
187class CheckVarsEscapingDeclContext final
188 : public ConstStmtVisitor<CheckVarsEscapingDeclContext> {
189 CodeGenFunction &CGF;
190 llvm::SetVector<const ValueDecl *> EscapedDecls;
191 llvm::SetVector<const ValueDecl *> EscapedVariableLengthDecls;
192 llvm::SetVector<const ValueDecl *> DelayedVariableLengthDecls;
193 llvm::SmallPtrSet<const Decl *, 4> EscapedParameters;
194 RecordDecl *GlobalizedRD = nullptr;
195 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
196 bool AllEscaped = false;
197 bool IsForCombinedParallelRegion = false;
198
199 void markAsEscaped(const ValueDecl *VD) {
200 // Do not globalize declare target variables.
201 if (!isa<VarDecl>(Val: VD) ||
202 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
203 return;
204 VD = cast<ValueDecl>(Val: VD->getCanonicalDecl());
205 // Use user-specified allocation.
206 if (VD->hasAttrs() && VD->hasAttr<OMPAllocateDeclAttr>())
207 return;
208 // Variables captured by value must be globalized.
209 bool IsCaptured = false;
210 if (auto *CSI = CGF.CapturedStmtInfo) {
211 if (const FieldDecl *FD = CSI->lookup(VD: cast<VarDecl>(Val: VD))) {
212 // Check if need to capture the variable that was already captured by
213 // value in the outer region.
214 IsCaptured = true;
215 if (!IsForCombinedParallelRegion) {
216 if (!FD->hasAttrs())
217 return;
218 const auto *Attr = FD->getAttr<OMPCaptureKindAttr>();
219 if (!Attr)
220 return;
221 if (((Attr->getCaptureKind() != OMPC_map) &&
222 !isOpenMPPrivate(Kind: Attr->getCaptureKind())) ||
223 ((Attr->getCaptureKind() == OMPC_map) &&
224 !FD->getType()->isAnyPointerType()))
225 return;
226 }
227 if (!FD->getType()->isReferenceType()) {
228 assert(!VD->getType()->isVariablyModifiedType() &&
229 "Parameter captured by value with variably modified type");
230 EscapedParameters.insert(Ptr: VD);
231 } else if (!IsForCombinedParallelRegion) {
232 return;
233 }
234 }
235 }
236 if ((!CGF.CapturedStmtInfo ||
237 (IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) &&
238 VD->getType()->isReferenceType())
239 // Do not globalize variables with reference type.
240 return;
241 if (VD->getType()->isVariablyModifiedType()) {
242 // If not captured at the target region level then mark the escaped
243 // variable as delayed.
244 if (IsCaptured)
245 EscapedVariableLengthDecls.insert(X: VD);
246 else
247 DelayedVariableLengthDecls.insert(X: VD);
248 } else
249 EscapedDecls.insert(X: VD);
250 }
251
252 void VisitValueDecl(const ValueDecl *VD) {
253 if (VD->getType()->isLValueReferenceType())
254 markAsEscaped(VD);
255 if (const auto *VarD = dyn_cast<VarDecl>(Val: VD)) {
256 if (!isa<ParmVarDecl>(Val: VarD) && VarD->hasInit()) {
257 const bool SavedAllEscaped = AllEscaped;
258 AllEscaped = VD->getType()->isLValueReferenceType();
259 Visit(S: VarD->getInit());
260 AllEscaped = SavedAllEscaped;
261 }
262 }
263 }
264 void VisitOpenMPCapturedStmt(const CapturedStmt *S,
265 ArrayRef<OMPClause *> Clauses,
266 bool IsCombinedParallelRegion) {
267 if (!S)
268 return;
269 for (const CapturedStmt::Capture &C : S->captures()) {
270 if (C.capturesVariable() && !C.capturesVariableByCopy()) {
271 const ValueDecl *VD = C.getCapturedVar();
272 bool SavedIsForCombinedParallelRegion = IsForCombinedParallelRegion;
273 if (IsCombinedParallelRegion) {
274 // Check if the variable is privatized in the combined construct and
275 // those private copies must be shared in the inner parallel
276 // directive.
277 IsForCombinedParallelRegion = false;
278 for (const OMPClause *C : Clauses) {
279 if (!isOpenMPPrivate(Kind: C->getClauseKind()) ||
280 C->getClauseKind() == OMPC_reduction ||
281 C->getClauseKind() == OMPC_linear ||
282 C->getClauseKind() == OMPC_private)
283 continue;
284 ArrayRef<const Expr *> Vars;
285 if (const auto *PC = dyn_cast<OMPFirstprivateClause>(Val: C))
286 Vars = PC->getVarRefs();
287 else if (const auto *PC = dyn_cast<OMPLastprivateClause>(Val: C))
288 Vars = PC->getVarRefs();
289 else
290 llvm_unreachable("Unexpected clause.");
291 for (const auto *E : Vars) {
292 const Decl *D =
293 cast<DeclRefExpr>(Val: E)->getDecl()->getCanonicalDecl();
294 if (D == VD->getCanonicalDecl()) {
295 IsForCombinedParallelRegion = true;
296 break;
297 }
298 }
299 if (IsForCombinedParallelRegion)
300 break;
301 }
302 }
303 markAsEscaped(VD);
304 if (isa<OMPCapturedExprDecl>(Val: VD))
305 VisitValueDecl(VD);
306 IsForCombinedParallelRegion = SavedIsForCombinedParallelRegion;
307 }
308 }
309 }
310
311 void buildRecordForGlobalizedVars(bool IsInTTDRegion) {
312 assert(!GlobalizedRD &&
313 "Record for globalized variables is built already.");
314 ArrayRef<const ValueDecl *> EscapedDeclsForParallel, EscapedDeclsForTeams;
315 unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;
316 if (IsInTTDRegion)
317 EscapedDeclsForTeams = EscapedDecls.getArrayRef();
318 else
319 EscapedDeclsForParallel = EscapedDecls.getArrayRef();
320 GlobalizedRD = ::buildRecordForGlobalizedVars(
321 C&: CGF.getContext(), EscapedDecls: EscapedDeclsForParallel, EscapedDeclsForTeams,
322 MappedDeclsFields, BufSize: WarpSize);
323 }
324
325public:
326 CheckVarsEscapingDeclContext(CodeGenFunction &CGF,
327 ArrayRef<const ValueDecl *> TeamsReductions)
328 : CGF(CGF), EscapedDecls(llvm::from_range, TeamsReductions) {}
329 ~CheckVarsEscapingDeclContext() = default;
330 void VisitDeclStmt(const DeclStmt *S) {
331 if (!S)
332 return;
333 for (const Decl *D : S->decls())
334 if (const auto *VD = dyn_cast_or_null<ValueDecl>(Val: D))
335 VisitValueDecl(VD);
336 }
337 void VisitOMPExecutableDirective(const OMPExecutableDirective *D) {
338 if (!D)
339 return;
340 if (!D->hasAssociatedStmt())
341 return;
342 if (const auto *S =
343 dyn_cast_or_null<CapturedStmt>(Val: D->getAssociatedStmt())) {
344 // Do not analyze directives that do not actually require capturing,
345 // like `omp for` or `omp simd` directives.
346 llvm::SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
347 getOpenMPCaptureRegions(CaptureRegions, DKind: D->getDirectiveKind());
348 if (CaptureRegions.size() == 1 && CaptureRegions.back() == OMPD_unknown) {
349 VisitStmt(S: S->getCapturedStmt());
350 return;
351 }
352 VisitOpenMPCapturedStmt(
353 S, Clauses: D->clauses(),
354 IsCombinedParallelRegion: CaptureRegions.back() == OMPD_parallel &&
355 isOpenMPDistributeDirective(DKind: D->getDirectiveKind()));
356 }
357 }
358 void VisitCapturedStmt(const CapturedStmt *S) {
359 if (!S)
360 return;
361 for (const CapturedStmt::Capture &C : S->captures()) {
362 if (C.capturesVariable() && !C.capturesVariableByCopy()) {
363 const ValueDecl *VD = C.getCapturedVar();
364 markAsEscaped(VD);
365 if (isa<OMPCapturedExprDecl>(Val: VD))
366 VisitValueDecl(VD);
367 }
368 }
369 }
370 void VisitLambdaExpr(const LambdaExpr *E) {
371 if (!E)
372 return;
373 for (const LambdaCapture &C : E->captures()) {
374 if (C.capturesVariable()) {
375 if (C.getCaptureKind() == LCK_ByRef) {
376 const ValueDecl *VD = C.getCapturedVar();
377 markAsEscaped(VD);
378 if (E->isInitCapture(Capture: &C) || isa<OMPCapturedExprDecl>(Val: VD))
379 VisitValueDecl(VD);
380 }
381 }
382 }
383 }
384 void VisitBlockExpr(const BlockExpr *E) {
385 if (!E)
386 return;
387 for (const BlockDecl::Capture &C : E->getBlockDecl()->captures()) {
388 if (C.isByRef()) {
389 const VarDecl *VD = C.getVariable();
390 markAsEscaped(VD);
391 if (isa<OMPCapturedExprDecl>(Val: VD) || VD->isInitCapture())
392 VisitValueDecl(VD);
393 }
394 }
395 }
396 void VisitCallExpr(const CallExpr *E) {
397 if (!E)
398 return;
399 for (const Expr *Arg : E->arguments()) {
400 if (!Arg)
401 continue;
402 if (Arg->isLValue()) {
403 const bool SavedAllEscaped = AllEscaped;
404 AllEscaped = true;
405 Visit(S: Arg);
406 AllEscaped = SavedAllEscaped;
407 } else {
408 Visit(S: Arg);
409 }
410 }
411 Visit(S: E->getCallee());
412 }
413 void VisitDeclRefExpr(const DeclRefExpr *E) {
414 if (!E)
415 return;
416 const ValueDecl *VD = E->getDecl();
417 if (AllEscaped)
418 markAsEscaped(VD);
419 if (isa<OMPCapturedExprDecl>(Val: VD))
420 VisitValueDecl(VD);
421 else if (VD->isInitCapture())
422 VisitValueDecl(VD);
423 }
424 void VisitUnaryOperator(const UnaryOperator *E) {
425 if (!E)
426 return;
427 if (E->getOpcode() == UO_AddrOf) {
428 const bool SavedAllEscaped = AllEscaped;
429 AllEscaped = true;
430 Visit(S: E->getSubExpr());
431 AllEscaped = SavedAllEscaped;
432 } else {
433 Visit(S: E->getSubExpr());
434 }
435 }
436 void VisitImplicitCastExpr(const ImplicitCastExpr *E) {
437 if (!E)
438 return;
439 if (E->getCastKind() == CK_ArrayToPointerDecay) {
440 const bool SavedAllEscaped = AllEscaped;
441 AllEscaped = true;
442 Visit(S: E->getSubExpr());
443 AllEscaped = SavedAllEscaped;
444 } else {
445 Visit(S: E->getSubExpr());
446 }
447 }
448 void VisitExpr(const Expr *E) {
449 if (!E)
450 return;
451 bool SavedAllEscaped = AllEscaped;
452 if (!E->isLValue())
453 AllEscaped = false;
454 for (const Stmt *Child : E->children())
455 if (Child)
456 Visit(S: Child);
457 AllEscaped = SavedAllEscaped;
458 }
459 void VisitStmt(const Stmt *S) {
460 if (!S)
461 return;
462 for (const Stmt *Child : S->children())
463 if (Child)
464 Visit(S: Child);
465 }
466
467 /// Returns the record that handles all the escaped local variables and used
468 /// instead of their original storage.
469 const RecordDecl *getGlobalizedRecord(bool IsInTTDRegion) {
470 if (!GlobalizedRD)
471 buildRecordForGlobalizedVars(IsInTTDRegion);
472 return GlobalizedRD;
473 }
474
475 /// Returns the field in the globalized record for the escaped variable.
476 const FieldDecl *getFieldForGlobalizedVar(const ValueDecl *VD) const {
477 assert(GlobalizedRD &&
478 "Record for globalized variables must be generated already.");
479 return MappedDeclsFields.lookup(Val: VD);
480 }
481
482 /// Returns the list of the escaped local variables/parameters.
483 ArrayRef<const ValueDecl *> getEscapedDecls() const {
484 return EscapedDecls.getArrayRef();
485 }
486
487 /// Checks if the escaped local variable is actually a parameter passed by
488 /// value.
489 const llvm::SmallPtrSetImpl<const Decl *> &getEscapedParameters() const {
490 return EscapedParameters;
491 }
492
493 /// Returns the list of the escaped variables with the variably modified
494 /// types.
495 ArrayRef<const ValueDecl *> getEscapedVariableLengthDecls() const {
496 return EscapedVariableLengthDecls.getArrayRef();
497 }
498
499 /// Returns the list of the delayed variables with the variably modified
500 /// types.
501 ArrayRef<const ValueDecl *> getDelayedVariableLengthDecls() const {
502 return DelayedVariableLengthDecls.getArrayRef();
503 }
504};
505} // anonymous namespace
506
507CGOpenMPRuntimeGPU::ExecutionMode
508CGOpenMPRuntimeGPU::getExecutionMode() const {
509 return CurrentExecutionMode;
510}
511
512CGOpenMPRuntimeGPU::DataSharingMode
513CGOpenMPRuntimeGPU::getDataSharingMode() const {
514 return CurrentDataSharingMode;
515}
516
517/// Check for inner (nested) SPMD construct, if any
518static bool hasNestedSPMDDirective(ASTContext &Ctx,
519 const OMPExecutableDirective &D) {
520 const auto *CS = D.getInnermostCapturedStmt();
521 const auto *Body =
522 CS->getCapturedStmt()->IgnoreContainers(/*IgnoreCaptured=*/true);
523 const Stmt *ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
524
525 if (const auto *NestedDir =
526 dyn_cast_or_null<OMPExecutableDirective>(Val: ChildStmt)) {
527 OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind();
528 switch (D.getDirectiveKind()) {
529 case OMPD_target:
530 if (isOpenMPParallelDirective(DKind))
531 return true;
532 if (DKind == OMPD_teams) {
533 Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
534 /*IgnoreCaptured=*/true);
535 if (!Body)
536 return false;
537 ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
538 if (const auto *NND =
539 dyn_cast_or_null<OMPExecutableDirective>(Val: ChildStmt)) {
540 DKind = NND->getDirectiveKind();
541 if (isOpenMPParallelDirective(DKind))
542 return true;
543 }
544 }
545 return false;
546 case OMPD_target_teams:
547 return isOpenMPParallelDirective(DKind);
548 case OMPD_target_simd:
549 case OMPD_target_parallel:
550 case OMPD_target_parallel_for:
551 case OMPD_target_parallel_for_simd:
552 case OMPD_target_teams_distribute:
553 case OMPD_target_teams_distribute_simd:
554 case OMPD_target_teams_distribute_parallel_for:
555 case OMPD_target_teams_distribute_parallel_for_simd:
556 case OMPD_parallel:
557 case OMPD_for:
558 case OMPD_parallel_for:
559 case OMPD_parallel_master:
560 case OMPD_parallel_sections:
561 case OMPD_for_simd:
562 case OMPD_parallel_for_simd:
563 case OMPD_cancel:
564 case OMPD_cancellation_point:
565 case OMPD_ordered_standalone:
566 case OMPD_ordered_blockassoc:
567 case OMPD_threadprivate:
568 case OMPD_allocate:
569 case OMPD_task:
570 case OMPD_simd:
571 case OMPD_sections:
572 case OMPD_section:
573 case OMPD_single:
574 case OMPD_master:
575 case OMPD_critical:
576 case OMPD_taskyield:
577 case OMPD_barrier:
578 case OMPD_taskwait:
579 case OMPD_taskgroup:
580 case OMPD_atomic:
581 case OMPD_flush:
582 case OMPD_depobj:
583 case OMPD_scan:
584 case OMPD_teams:
585 case OMPD_target_data:
586 case OMPD_target_exit_data:
587 case OMPD_target_enter_data:
588 case OMPD_distribute:
589 case OMPD_distribute_simd:
590 case OMPD_distribute_parallel_for:
591 case OMPD_distribute_parallel_for_simd:
592 case OMPD_teams_distribute:
593 case OMPD_teams_distribute_simd:
594 case OMPD_teams_distribute_parallel_for:
595 case OMPD_teams_distribute_parallel_for_simd:
596 case OMPD_target_update:
597 case OMPD_declare_simd:
598 case OMPD_declare_variant:
599 case OMPD_begin_declare_variant:
600 case OMPD_end_declare_variant:
601 case OMPD_declare_target:
602 case OMPD_end_declare_target:
603 case OMPD_declare_reduction:
604 case OMPD_declare_mapper:
605 case OMPD_taskloop:
606 case OMPD_taskloop_simd:
607 case OMPD_master_taskloop:
608 case OMPD_master_taskloop_simd:
609 case OMPD_parallel_master_taskloop:
610 case OMPD_parallel_master_taskloop_simd:
611 case OMPD_requires:
612 case OMPD_unknown:
613 default:
614 llvm_unreachable("Unexpected directive.");
615 }
616 }
617
618 return false;
619}
620
621static bool supportsSPMDExecutionMode(ASTContext &Ctx,
622 const OMPExecutableDirective &D) {
623 OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
624 switch (DirectiveKind) {
625 case OMPD_target:
626 case OMPD_target_teams:
627 return hasNestedSPMDDirective(Ctx, D);
628 case OMPD_target_parallel_loop:
629 case OMPD_target_parallel:
630 case OMPD_target_parallel_for:
631 case OMPD_target_parallel_for_simd:
632 case OMPD_target_teams_distribute_parallel_for:
633 case OMPD_target_teams_distribute_parallel_for_simd:
634 case OMPD_target_simd:
635 case OMPD_target_teams_distribute_simd:
636 return true;
637 case OMPD_target_teams_distribute:
638 return false;
639 case OMPD_target_teams_loop:
640 // Whether this is true or not depends on how the directive will
641 // eventually be emitted.
642 if (auto *TTLD = dyn_cast<OMPTargetTeamsGenericLoopDirective>(Val: &D))
643 return TTLD->canBeParallelFor();
644 return false;
645 case OMPD_parallel:
646 case OMPD_for:
647 case OMPD_parallel_for:
648 case OMPD_parallel_master:
649 case OMPD_parallel_sections:
650 case OMPD_for_simd:
651 case OMPD_parallel_for_simd:
652 case OMPD_cancel:
653 case OMPD_cancellation_point:
654 case OMPD_ordered_standalone:
655 case OMPD_ordered_blockassoc:
656 case OMPD_threadprivate:
657 case OMPD_allocate:
658 case OMPD_task:
659 case OMPD_simd:
660 case OMPD_sections:
661 case OMPD_section:
662 case OMPD_single:
663 case OMPD_master:
664 case OMPD_critical:
665 case OMPD_taskyield:
666 case OMPD_barrier:
667 case OMPD_taskwait:
668 case OMPD_taskgroup:
669 case OMPD_atomic:
670 case OMPD_flush:
671 case OMPD_depobj:
672 case OMPD_scan:
673 case OMPD_teams:
674 case OMPD_target_data:
675 case OMPD_target_exit_data:
676 case OMPD_target_enter_data:
677 case OMPD_distribute:
678 case OMPD_distribute_simd:
679 case OMPD_distribute_parallel_for:
680 case OMPD_distribute_parallel_for_simd:
681 case OMPD_teams_distribute:
682 case OMPD_teams_distribute_simd:
683 case OMPD_teams_distribute_parallel_for:
684 case OMPD_teams_distribute_parallel_for_simd:
685 case OMPD_target_update:
686 case OMPD_declare_simd:
687 case OMPD_declare_variant:
688 case OMPD_begin_declare_variant:
689 case OMPD_end_declare_variant:
690 case OMPD_declare_target:
691 case OMPD_end_declare_target:
692 case OMPD_declare_reduction:
693 case OMPD_declare_mapper:
694 case OMPD_taskloop:
695 case OMPD_taskloop_simd:
696 case OMPD_master_taskloop:
697 case OMPD_master_taskloop_simd:
698 case OMPD_parallel_master_taskloop:
699 case OMPD_parallel_master_taskloop_simd:
700 case OMPD_requires:
701 case OMPD_unknown:
702 default:
703 break;
704 }
705 llvm_unreachable(
706 "Unknown programming model for OpenMP directive on NVPTX target.");
707}
708
709void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
710 StringRef ParentName,
711 llvm::Function *&OutlinedFn,
712 llvm::Constant *&OutlinedFnID,
713 bool IsOffloadEntry,
714 const RegionCodeGenTy &CodeGen) {
715 ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_NonSPMD);
716 EntryFunctionState EST;
717 WrapperFunctionsMap.clear();
718
719 [[maybe_unused]] bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
720 assert(!IsBareKernel && "bare kernel should not be at generic mode");
721
722 // Emit target region as a standalone region.
723 class NVPTXPrePostActionTy : public PrePostActionTy {
724 CGOpenMPRuntimeGPU::EntryFunctionState &EST;
725 const OMPExecutableDirective &D;
726
727 public:
728 NVPTXPrePostActionTy(CGOpenMPRuntimeGPU::EntryFunctionState &EST,
729 const OMPExecutableDirective &D)
730 : EST(EST), D(D) {}
731 void Enter(CodeGenFunction &CGF) override {
732 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
733 RT.emitKernelInit(D, CGF, EST, /* IsSPMD */ false);
734 // Skip target region initialization.
735 RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
736 }
737 void Exit(CodeGenFunction &CGF) override {
738 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
739 RT.clearLocThreadIdInsertPt(CGF);
740 RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ false);
741 }
742 } Action(EST, D);
743 CodeGen.setAction(Action);
744 IsInTTDRegion = true;
745 emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
746 IsOffloadEntry, CodeGen);
747 IsInTTDRegion = false;
748}
749
750void CGOpenMPRuntimeGPU::emitBareKernelEnvironment(
751 const OMPExecutableDirective &D, CodeGenFunction &CGF) {
752 // Bare kernels manage their own initialization and never call
753 // __kmpc_target_init, but the runtime still needs a
754 // '<kernel>_kernel_environment' global to know how the kernel was
755 // configured, so emit it directly here.
756 llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;
757 Attrs.ExecFlags = llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_BARE;
758 CGBuilderTy &Bld = CGF.Builder;
759 OMPBuilder.emitKernelEnvironment(Loc: Bld, Attrs);
760}
761
762void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D,
763 CodeGenFunction &CGF,
764 EntryFunctionState &EST, bool IsSPMD) {
765 llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;
766 Attrs.ExecFlags =
767 IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD
768 : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;
769 computeMinAndMaxThreadsAndTeams(D, CGF, Attrs);
770
771 CGBuilderTy &Bld = CGF.Builder;
772 Bld.restoreIP(IP: OMPBuilder.createTargetInit(Loc: Bld, Attrs));
773 if (!IsSPMD)
774 emitGenericVarsProlog(CGF, Loc: EST.Loc);
775}
776
777void CGOpenMPRuntimeGPU::emitKernelDeinit(CodeGenFunction &CGF,
778 EntryFunctionState &EST,
779 bool IsSPMD) {
780 if (!IsSPMD)
781 emitGenericVarsEpilog(CGF);
782
783 // This is temporary until we remove the fixed sized buffer.
784 ASTContext &C = CGM.getContext();
785 RecordDecl *StaticRD = C.buildImplicitRecord(
786 Name: "_openmp_teams_reduction_type_$_", TK: RecordDecl::TagKind::Union);
787 StaticRD->startDefinition();
788 for (const RecordDecl *TeamReductionRec : TeamsReductions) {
789 CanQualType RecTy = C.getCanonicalTagType(TD: TeamReductionRec);
790 auto *Field = FieldDecl::Create(
791 C, DC: StaticRD, StartLoc: SourceLocation(), IdLoc: SourceLocation(), Id: nullptr, T: RecTy,
792 TInfo: C.getTrivialTypeSourceInfo(T: RecTy, Loc: SourceLocation()),
793 /*BW=*/nullptr, /*Mutable=*/false,
794 /*InitStyle=*/ICIS_NoInit);
795 Field->setAccess(AS_public);
796 StaticRD->addDecl(D: Field);
797 }
798 StaticRD->completeDefinition();
799 CanQualType StaticTy = C.getCanonicalTagType(TD: StaticRD);
800 llvm::Type *LLVMReductionsBufferTy =
801 CGM.getTypes().ConvertTypeForMem(T: StaticTy);
802 const auto &DL = CGM.getModule().getDataLayout();
803 uint64_t ReductionDataSize =
804 TeamsReductions.empty()
805 ? 0
806 : DL.getTypeAllocSize(Ty: LLVMReductionsBufferTy).getFixedValue();
807 CGBuilderTy &Bld = CGF.Builder;
808 OMPBuilder.createTargetDeinit(Loc: Bld, TeamsReductionDataSize: ReductionDataSize);
809 TeamsReductions.clear();
810}
811
812void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
813 StringRef ParentName,
814 llvm::Function *&OutlinedFn,
815 llvm::Constant *&OutlinedFnID,
816 bool IsOffloadEntry,
817 const RegionCodeGenTy &CodeGen) {
818 ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_SPMD);
819 EntryFunctionState EST;
820
821 bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
822
823 // Emit target region as a standalone region.
824 class NVPTXPrePostActionTy : public PrePostActionTy {
825 CGOpenMPRuntimeGPU &RT;
826 CGOpenMPRuntimeGPU::EntryFunctionState &EST;
827 bool IsBareKernel;
828 DataSharingMode Mode;
829 const OMPExecutableDirective &D;
830
831 public:
832 NVPTXPrePostActionTy(CGOpenMPRuntimeGPU &RT,
833 CGOpenMPRuntimeGPU::EntryFunctionState &EST,
834 bool IsBareKernel, const OMPExecutableDirective &D)
835 : RT(RT), EST(EST), IsBareKernel(IsBareKernel),
836 Mode(RT.CurrentDataSharingMode), D(D) {}
837 void Enter(CodeGenFunction &CGF) override {
838 if (IsBareKernel) {
839 RT.CurrentDataSharingMode = DataSharingMode::DS_CUDA;
840 RT.emitBareKernelEnvironment(D, CGF);
841 return;
842 }
843 RT.emitKernelInit(D, CGF, EST, /* IsSPMD */ true);
844 // Skip target region initialization.
845 RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
846 }
847 void Exit(CodeGenFunction &CGF) override {
848 if (IsBareKernel) {
849 RT.CurrentDataSharingMode = Mode;
850 return;
851 }
852 RT.clearLocThreadIdInsertPt(CGF);
853 RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ true);
854 }
855 } Action(*this, EST, IsBareKernel, D);
856 CodeGen.setAction(Action);
857 IsInTTDRegion = true;
858 emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
859 IsOffloadEntry, CodeGen);
860 IsInTTDRegion = false;
861}
862
863void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
864 const OMPExecutableDirective &D, StringRef ParentName,
865 llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
866 bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
867 if (!IsOffloadEntry) // Nothing to do.
868 return;
869
870 assert(!ParentName.empty() && "Invalid target region parent name!");
871
872 bool Mode = supportsSPMDExecutionMode(Ctx&: CGM.getContext(), D);
873 bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
874 if (Mode || IsBareKernel)
875 emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
876 CodeGen);
877 else
878 emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
879 CodeGen);
880}
881
882CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
883 : CGOpenMPRuntime(CGM) {
884 llvm::OpenMPIRBuilderConfig Config(
885 CGM.getLangOpts().OpenMPIsTargetDevice, isGPU(),
886 CGM.getLangOpts().OpenMPOffloadMandatory,
887 /*HasRequiresReverseOffload*/ false, /*HasRequiresUnifiedAddress*/ false,
888 hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
889 Config.setDefaultTargetAS(
890 CGM.getContext().getTargetInfo().getTargetAddressSpace(AS: LangAS::Default));
891 Config.setRuntimeCC(CGM.getRuntimeCC());
892
893 OMPBuilder.setConfig(Config);
894
895 if (!CGM.getLangOpts().OpenMPIsTargetDevice)
896 llvm_unreachable("OpenMP can only handle device code.");
897
898 if (CGM.getLangOpts().OpenMPCUDAMode)
899 CurrentDataSharingMode = CGOpenMPRuntimeGPU::DS_CUDA;
900
901 llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder();
902 if (CGM.getLangOpts().NoGPULib || CGM.getLangOpts().OMPHostIRFile.empty())
903 return;
904
905 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPTargetDebug,
906 Name: "__omp_rtl_debug_kind");
907 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPTeamSubscription,
908 Name: "__omp_rtl_assume_teams_oversubscription");
909 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPThreadSubscription,
910 Name: "__omp_rtl_assume_threads_oversubscription");
911 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPNoThreadState,
912 Name: "__omp_rtl_assume_no_thread_state");
913 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPNoNestedParallelism,
914 Name: "__omp_rtl_assume_no_nested_parallelism");
915}
916
917void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,
918 ProcBindKind ProcBind,
919 SourceLocation Loc) {
920 // Nothing to do.
921}
922
923llvm::Value *CGOpenMPRuntimeGPU::emitMessageClause(CodeGenFunction &CGF,
924 const Expr *Message,
925 SourceLocation Loc) {
926 CGM.getDiags().Report(Loc, DiagID: diag::warn_omp_gpu_unsupported_clause)
927 << getOpenMPClauseName(C: OMPC_message);
928 return nullptr;
929}
930
931llvm::Value *
932CGOpenMPRuntimeGPU::emitSeverityClause(OpenMPSeverityClauseKind Severity,
933 SourceLocation Loc) {
934 CGM.getDiags().Report(Loc, DiagID: diag::warn_omp_gpu_unsupported_clause)
935 << getOpenMPClauseName(C: OMPC_severity);
936 return nullptr;
937}
938
939void CGOpenMPRuntimeGPU::emitNumThreadsClause(
940 CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
941 OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,
942 SourceLocation SeverityLoc, const Expr *Message,
943 SourceLocation MessageLoc) {
944 if (Modifier == OMPC_NUMTHREADS_strict) {
945 CGM.getDiags().Report(Loc,
946 DiagID: diag::warn_omp_gpu_unsupported_modifier_for_clause)
947 << "strict" << getOpenMPClauseName(C: OMPC_num_threads);
948 return;
949 }
950
951 // Nothing to do.
952}
953
954void CGOpenMPRuntimeGPU::emitNumTeamsClause(CodeGenFunction &CGF,
955 const Expr *NumTeams,
956 const Expr *ThreadLimit,
957 SourceLocation Loc) {}
958
959llvm::Function *CGOpenMPRuntimeGPU::emitParallelOutlinedFunction(
960 CodeGenFunction &CGF, const OMPExecutableDirective &D,
961 const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
962 const RegionCodeGenTy &CodeGen) {
963 // Emit target region as a standalone region.
964 bool PrevIsInTTDRegion = IsInTTDRegion;
965 IsInTTDRegion = false;
966 auto *OutlinedFun =
967 cast<llvm::Function>(Val: CGOpenMPRuntime::emitParallelOutlinedFunction(
968 CGF, D, ThreadIDVar, InnermostKind, CodeGen));
969 IsInTTDRegion = PrevIsInTTDRegion;
970 if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD) {
971 llvm::Function *WrapperFun =
972 createParallelDataSharingWrapper(OutlinedParallelFn: OutlinedFun, D);
973 WrapperFunctionsMap[OutlinedFun] = WrapperFun;
974 }
975
976 return OutlinedFun;
977}
978
979/// Get list of lastprivate variables from the teams distribute ... or
980/// teams {distribute ...} directives.
981static void
982getDistributeLastprivateVars(ASTContext &Ctx, const OMPExecutableDirective &D,
983 llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
984 assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
985 "expected teams directive.");
986 const OMPExecutableDirective *Dir = &D;
987 if (!isOpenMPDistributeDirective(DKind: D.getDirectiveKind())) {
988 if (const Stmt *S = CGOpenMPRuntime::getSingleCompoundChild(
989 Ctx,
990 Body: D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(
991 /*IgnoreCaptured=*/true))) {
992 Dir = dyn_cast_or_null<OMPExecutableDirective>(Val: S);
993 if (Dir && !isOpenMPDistributeDirective(DKind: Dir->getDirectiveKind()))
994 Dir = nullptr;
995 }
996 }
997 if (!Dir)
998 return;
999 for (const auto *C : Dir->getClausesOfKind<OMPLastprivateClause>()) {
1000 for (const Expr *E : C->getVarRefs())
1001 Vars.push_back(Elt: getPrivateItem(RefExpr: E));
1002 }
1003}
1004
1005/// Get list of reduction variables from the teams ... directives.
1006static void
1007getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D,
1008 llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
1009 assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
1010 "expected teams directive.");
1011 for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
1012 for (const Expr *E : C->privates())
1013 Vars.push_back(Elt: getPrivateItem(RefExpr: E));
1014 }
1015}
1016
1017llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(
1018 CodeGenFunction &CGF, const OMPExecutableDirective &D,
1019 const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
1020 const RegionCodeGenTy &CodeGen) {
1021 SourceLocation Loc = D.getBeginLoc();
1022
1023 const RecordDecl *GlobalizedRD = nullptr;
1024 llvm::SmallVector<const ValueDecl *, 4> LastPrivatesReductions;
1025 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
1026 unsigned WarpSize = CGM.getTarget().getGridValue().GV_Warp_Size;
1027 // Globalize team reductions variable unconditionally in all modes.
1028 if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1029 getTeamsReductionVars(Ctx&: CGM.getContext(), D, Vars&: LastPrivatesReductions);
1030 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
1031 getDistributeLastprivateVars(Ctx&: CGM.getContext(), D, Vars&: LastPrivatesReductions);
1032 if (!LastPrivatesReductions.empty()) {
1033 GlobalizedRD = ::buildRecordForGlobalizedVars(
1034 C&: CGM.getContext(), EscapedDecls: {}, EscapedDeclsForTeams: LastPrivatesReductions, MappedDeclsFields,
1035 BufSize: WarpSize);
1036 }
1037 } else if (!LastPrivatesReductions.empty()) {
1038 assert(!TeamAndReductions.first &&
1039 "Previous team declaration is not expected.");
1040 TeamAndReductions.first = D.getCapturedStmt(RegionKind: OMPD_teams)->getCapturedDecl();
1041 std::swap(LHS&: TeamAndReductions.second, RHS&: LastPrivatesReductions);
1042 }
1043
1044 // Emit target region as a standalone region.
1045 class NVPTXPrePostActionTy : public PrePostActionTy {
1046 SourceLocation &Loc;
1047 const RecordDecl *GlobalizedRD;
1048 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1049 &MappedDeclsFields;
1050
1051 public:
1052 NVPTXPrePostActionTy(
1053 SourceLocation &Loc, const RecordDecl *GlobalizedRD,
1054 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1055 &MappedDeclsFields)
1056 : Loc(Loc), GlobalizedRD(GlobalizedRD),
1057 MappedDeclsFields(MappedDeclsFields) {}
1058 void Enter(CodeGenFunction &CGF) override {
1059 auto &Rt =
1060 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1061 if (GlobalizedRD) {
1062 auto I = Rt.FunctionGlobalizedDecls.try_emplace(Key: CGF.CurFn).first;
1063 I->getSecond().MappedParams =
1064 std::make_unique<CodeGenFunction::OMPMapVars>();
1065 DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
1066 for (const auto &Pair : MappedDeclsFields) {
1067 assert(Pair.getFirst()->isCanonicalDecl() &&
1068 "Expected canonical declaration");
1069 Data.try_emplace(Key: Pair.getFirst());
1070 }
1071 }
1072 Rt.emitGenericVarsProlog(CGF, Loc);
1073 }
1074 void Exit(CodeGenFunction &CGF) override {
1075 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
1076 .emitGenericVarsEpilog(CGF);
1077 }
1078 } Action(Loc, GlobalizedRD, MappedDeclsFields);
1079 CodeGen.setAction(Action);
1080 llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(
1081 CGF, D, ThreadIDVar, InnermostKind, CodeGen);
1082
1083 return OutlinedFun;
1084}
1085
1086void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,
1087 SourceLocation Loc) {
1088 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
1089 return;
1090
1091 CGBuilderTy &Bld = CGF.Builder;
1092
1093 const auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
1094 if (I == FunctionGlobalizedDecls.end())
1095 return;
1096
1097 for (auto &Rec : I->getSecond().LocalVarData) {
1098 const auto *VD = cast<VarDecl>(Val: Rec.first);
1099 bool EscapedParam = I->getSecond().EscapedParameters.count(Ptr: Rec.first);
1100 QualType VarTy = VD->getType();
1101
1102 // Get the local allocation of a firstprivate variable before sharing
1103 llvm::Value *ParValue;
1104 if (EscapedParam) {
1105 LValue ParLVal =
1106 CGF.MakeAddrLValue(Addr: CGF.GetAddrOfLocalVar(VD), T: VD->getType());
1107 ParValue = CGF.EmitLoadOfScalar(lvalue: ParLVal, Loc);
1108 }
1109
1110 // Allocate space for the variable to be globalized
1111 llvm::Value *AllocArgs[] = {CGF.getTypeSize(Ty: VD->getType())};
1112 llvm::CallBase *VoidPtr =
1113 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1114 M&: CGM.getModule(), FnID: OMPRTL___kmpc_alloc_shared),
1115 args: AllocArgs, name: VD->getName());
1116 // FIXME: We should use the variables actual alignment as an argument.
1117 VoidPtr->addRetAttr(Attr: llvm::Attribute::get(
1118 Context&: CGM.getLLVMContext(), Kind: llvm::Attribute::Alignment,
1119 Val: CGM.getContext().getTargetInfo().getNewAlign() / 8));
1120
1121 // Cast the void pointer and get the address of the globalized variable.
1122 llvm::Value *CastedVoidPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1123 V: VoidPtr, DestTy: Bld.getPtrTy(AddrSpace: 0), Name: VD->getName() + "_on_stack");
1124 LValue VarAddr =
1125 CGF.MakeNaturalAlignPointeeRawAddrLValue(V: CastedVoidPtr, T: VarTy);
1126 Rec.second.PrivateAddr = VarAddr.getAddress();
1127 Rec.second.GlobalizedVal = VoidPtr;
1128
1129 // Assign the local allocation to the newly globalized location.
1130 if (EscapedParam) {
1131 CGF.EmitStoreOfScalar(value: ParValue, lvalue: VarAddr);
1132 I->getSecond().MappedParams->setVarAddr(CGF, LocalVD: VD, TempAddr: VarAddr.getAddress());
1133 }
1134 if (auto *DI = CGF.getDebugInfo())
1135 VoidPtr->setDebugLoc(DI->SourceLocToDebugLoc(Loc: VD->getLocation()));
1136 }
1137
1138 for (const auto *ValueD : I->getSecond().EscapedVariableLengthDecls) {
1139 const auto *VD = cast<VarDecl>(Val: ValueD);
1140 std::pair<llvm::Value *, llvm::Value *> AddrSizePair =
1141 getKmpcAllocShared(CGF, VD);
1142 I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(Args&: AddrSizePair);
1143 LValue Base = CGF.MakeAddrLValue(V: AddrSizePair.first, T: VD->getType(),
1144 Alignment: CGM.getContext().getDeclAlign(D: VD),
1145 Source: AlignmentSource::Decl);
1146 I->getSecond().MappedParams->setVarAddr(CGF, LocalVD: VD, TempAddr: Base.getAddress());
1147 }
1148 I->getSecond().MappedParams->apply(CGF);
1149}
1150
1151bool CGOpenMPRuntimeGPU::isDelayedVariableLengthDecl(CodeGenFunction &CGF,
1152 const VarDecl *VD) const {
1153 const auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
1154 if (I == FunctionGlobalizedDecls.end())
1155 return false;
1156
1157 // Check variable declaration is delayed:
1158 return llvm::is_contained(Range: I->getSecond().DelayedVariableLengthDecls, Element: VD);
1159}
1160
1161std::pair<llvm::Value *, llvm::Value *>
1162CGOpenMPRuntimeGPU::getKmpcAllocShared(CodeGenFunction &CGF,
1163 const VarDecl *VD) {
1164 CGBuilderTy &Bld = CGF.Builder;
1165
1166 // Compute size and alignment.
1167 llvm::Value *Size = CGF.getTypeSize(Ty: VD->getType());
1168 CharUnits Align = CGM.getContext().getDeclAlign(D: VD);
1169 Size = Bld.CreateNUWAdd(
1170 LHS: Size, RHS: llvm::ConstantInt::get(Ty: CGF.SizeTy, V: Align.getQuantity() - 1));
1171 llvm::Value *AlignVal =
1172 llvm::ConstantInt::get(Ty: CGF.SizeTy, V: Align.getQuantity());
1173 Size = Bld.CreateUDiv(LHS: Size, RHS: AlignVal);
1174 Size = Bld.CreateNUWMul(LHS: Size, RHS: AlignVal);
1175
1176 // Allocate space for this VLA object to be globalized.
1177 llvm::Value *AllocArgs[] = {Size};
1178 llvm::CallBase *VoidPtr =
1179 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1180 M&: CGM.getModule(), FnID: OMPRTL___kmpc_alloc_shared),
1181 args: AllocArgs, name: VD->getName());
1182 VoidPtr->addRetAttr(Attr: llvm::Attribute::get(
1183 Context&: CGM.getLLVMContext(), Kind: llvm::Attribute::Alignment, Val: Align.getQuantity()));
1184
1185 return std::make_pair(x&: VoidPtr, y&: Size);
1186}
1187
1188void CGOpenMPRuntimeGPU::getKmpcFreeShared(
1189 CodeGenFunction &CGF,
1190 const std::pair<llvm::Value *, llvm::Value *> &AddrSizePair) {
1191 // Deallocate the memory for each globalized VLA object
1192 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1193 M&: CGM.getModule(), FnID: OMPRTL___kmpc_free_shared),
1194 args: {AddrSizePair.first, AddrSizePair.second});
1195}
1196
1197void CGOpenMPRuntimeGPU::emitGenericVarsEpilog(CodeGenFunction &CGF) {
1198 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
1199 return;
1200
1201 const auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
1202 if (I != FunctionGlobalizedDecls.end()) {
1203 // Deallocate the memory for each globalized VLA object that was
1204 // globalized in the prolog (i.e. emitGenericVarsProlog).
1205 for (const auto &AddrSizePair :
1206 llvm::reverse(C&: I->getSecond().EscapedVariableLengthDeclsAddrs)) {
1207 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1208 M&: CGM.getModule(), FnID: OMPRTL___kmpc_free_shared),
1209 args: {AddrSizePair.first, AddrSizePair.second});
1210 }
1211 // Deallocate the memory for each globalized value
1212 for (auto &Rec : llvm::reverse(C&: I->getSecond().LocalVarData)) {
1213 const auto *VD = cast<VarDecl>(Val: Rec.first);
1214 I->getSecond().MappedParams->restore(CGF);
1215
1216 llvm::Value *FreeArgs[] = {Rec.second.GlobalizedVal,
1217 CGF.getTypeSize(Ty: VD->getType())};
1218 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1219 M&: CGM.getModule(), FnID: OMPRTL___kmpc_free_shared),
1220 args: FreeArgs);
1221 }
1222 }
1223}
1224
1225void CGOpenMPRuntimeGPU::emitTeamsCall(CodeGenFunction &CGF,
1226 const OMPExecutableDirective &D,
1227 SourceLocation Loc,
1228 llvm::Function *OutlinedFn,
1229 ArrayRef<llvm::Value *> CapturedVars) {
1230 if (!CGF.HaveInsertPoint())
1231 return;
1232
1233 bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
1234
1235 RawAddress ZeroAddr = CGF.CreateDefaultAlignTempAlloca(Ty: CGF.Int32Ty,
1236 /*Name=*/".zero.addr");
1237 CGF.Builder.CreateStore(Val: CGF.Builder.getInt32(/*C*/ 0), Addr: ZeroAddr);
1238 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
1239 // We don't emit any thread id function call in bare kernel, but because the
1240 // outlined function has a pointer argument, we emit a nullptr here.
1241 if (IsBareKernel)
1242 OutlinedFnArgs.push_back(Elt: llvm::ConstantPointerNull::get(T: CGM.VoidPtrTy));
1243 else
1244 OutlinedFnArgs.push_back(Elt: emitThreadIDAddress(CGF, Loc).emitRawPointer(CGF));
1245 OutlinedFnArgs.push_back(Elt: ZeroAddr.getPointer());
1246 OutlinedFnArgs.append(in_start: CapturedVars.begin(), in_end: CapturedVars.end());
1247 emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, Args: OutlinedFnArgs);
1248}
1249
1250void CGOpenMPRuntimeGPU::emitParallelCall(
1251 CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *OutlinedFn,
1252 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond,
1253 llvm::Value *NumThreads, OpenMPNumThreadsClauseModifier NumThreadsModifier,
1254 OpenMPSeverityClauseKind Severity, const Expr *Message) {
1255 if (!CGF.HaveInsertPoint())
1256 return;
1257
1258 auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
1259 NumThreads](CodeGenFunction &CGF,
1260 PrePostActionTy &Action) {
1261 CGBuilderTy &Bld = CGF.Builder;
1262 llvm::Value *NumThreadsVal = NumThreads;
1263 llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
1264 llvm::PointerType *FnPtrTy = llvm::PointerType::get(
1265 C&: CGF.getLLVMContext(), AddressSpace: CGM.getDataLayout().getProgramAddressSpace());
1266
1267 llvm::Value *ID = llvm::ConstantPointerNull::get(T: FnPtrTy);
1268 if (WFn)
1269 ID = Bld.CreateBitOrPointerCast(V: WFn, DestTy: FnPtrTy);
1270
1271 llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(V: OutlinedFn, DestTy: FnPtrTy);
1272
1273 // Create a private scope that will globalize the arguments
1274 // passed from the outside of the target region.
1275 // TODO: Is that needed?
1276 CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF);
1277
1278 Address CapturedVarsAddrs = CGF.CreateDefaultAlignTempAlloca(
1279 Ty: llvm::ArrayType::get(ElementType: CGM.VoidPtrTy, NumElements: CapturedVars.size()),
1280 Name: "captured_vars_addrs");
1281 // There's something to share.
1282 if (!CapturedVars.empty()) {
1283 // Prepare for parallel region. Indicate the outlined function.
1284 ASTContext &Ctx = CGF.getContext();
1285 unsigned Idx = 0;
1286 for (llvm::Value *V : CapturedVars) {
1287 Address Dst = Bld.CreateConstArrayGEP(Addr: CapturedVarsAddrs, Index: Idx);
1288 llvm::Value *PtrV;
1289 if (V->getType()->isIntegerTy())
1290 PtrV = Bld.CreateIntToPtr(V, DestTy: CGF.VoidPtrTy);
1291 else
1292 PtrV = Bld.CreatePointerBitCastOrAddrSpaceCast(V, DestTy: CGF.VoidPtrTy);
1293 CGF.EmitStoreOfScalar(Value: PtrV, Addr: Dst, /*Volatile=*/false,
1294 Ty: Ctx.getPointerType(T: Ctx.VoidPtrTy));
1295 ++Idx;
1296 }
1297 }
1298
1299 llvm::Value *IfCondVal = nullptr;
1300 if (IfCond)
1301 IfCondVal = Bld.CreateIntCast(V: CGF.EvaluateExprAsBool(E: IfCond), DestTy: CGF.Int32Ty,
1302 /* isSigned */ false);
1303 else
1304 IfCondVal = llvm::ConstantInt::get(Ty: CGF.Int32Ty, V: 1);
1305
1306 if (!NumThreadsVal)
1307 NumThreadsVal = llvm::ConstantInt::getAllOnesValue(Ty: CGF.Int32Ty);
1308 else
1309 NumThreadsVal = Bld.CreateZExtOrTrunc(V: NumThreadsVal, DestTy: CGF.Int32Ty);
1310
1311 // No strict prescriptiveness for the number of threads.
1312 llvm::Value *StrictNumThreadsVal = llvm::ConstantInt::get(Ty: CGF.Int32Ty, V: 0);
1313
1314 assert(IfCondVal && "Expected a value");
1315 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
1316 llvm::Value *Args[] = {
1317 RTLoc,
1318 getThreadID(CGF, Loc),
1319 IfCondVal,
1320 NumThreadsVal,
1321 llvm::ConstantInt::getAllOnesValue(Ty: CGF.Int32Ty),
1322 FnPtr,
1323 ID,
1324 Bld.CreateBitOrPointerCast(V: CapturedVarsAddrs.emitRawPointer(CGF),
1325 DestTy: CGF.VoidPtrPtrTy),
1326 llvm::ConstantInt::get(Ty: CGM.SizeTy, V: CapturedVars.size()),
1327 StrictNumThreadsVal};
1328
1329 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1330 M&: CGM.getModule(), FnID: OMPRTL___kmpc_parallel_60),
1331 args: Args);
1332 };
1333
1334 RegionCodeGenTy RCG(ParallelGen);
1335 RCG(CGF);
1336}
1337
1338void CGOpenMPRuntimeGPU::syncCTAThreads(CodeGenFunction &CGF) {
1339 // Always emit simple barriers!
1340 if (!CGF.HaveInsertPoint())
1341 return;
1342 // Build call __kmpc_barrier_simple_spmd(nullptr, 0);
1343 // This function does not use parameters, so we can emit just default values.
1344 llvm::Value *Args[] = {
1345 llvm::ConstantPointerNull::get(
1346 T: cast<llvm::PointerType>(Val: getIdentTyPointerTy())),
1347 llvm::ConstantInt::get(Ty: CGF.Int32Ty, /*V=*/0, /*isSigned=*/IsSigned: true)};
1348 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1349 M&: CGM.getModule(), FnID: OMPRTL___kmpc_barrier_simple_spmd),
1350 args: Args);
1351}
1352
1353void CGOpenMPRuntimeGPU::emitBarrierCall(CodeGenFunction &CGF,
1354 SourceLocation Loc,
1355 OpenMPDirectiveKind Kind, bool,
1356 bool) {
1357 // Always emit simple barriers!
1358 if (!CGF.HaveInsertPoint())
1359 return;
1360 // Build call __kmpc_cancel_barrier(loc, thread_id);
1361 unsigned Flags = getDefaultFlagsForBarriers(Kind);
1362 llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc, Flags),
1363 getThreadID(CGF, Loc)};
1364
1365 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1366 M&: CGM.getModule(), FnID: OMPRTL___kmpc_barrier),
1367 args: Args);
1368}
1369
1370void CGOpenMPRuntimeGPU::emitCriticalRegion(
1371 CodeGenFunction &CGF, StringRef CriticalName,
1372 const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc,
1373 const Expr *Hint) {
1374 llvm::BasicBlock *LoopBB = CGF.createBasicBlock(name: "omp.critical.loop");
1375 llvm::BasicBlock *TestBB = CGF.createBasicBlock(name: "omp.critical.test");
1376 llvm::BasicBlock *SyncBB = CGF.createBasicBlock(name: "omp.critical.sync");
1377 llvm::BasicBlock *BodyBB = CGF.createBasicBlock(name: "omp.critical.body");
1378 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(name: "omp.critical.exit");
1379
1380 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1381
1382 // Get the mask of active threads in the warp.
1383 llvm::Value *Mask = CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1384 M&: CGM.getModule(), FnID: OMPRTL___kmpc_warp_active_thread_mask));
1385 // Fetch team-local id of the thread.
1386 llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
1387
1388 // Get the width of the team.
1389 llvm::Value *TeamWidth = RT.getGPUNumThreads(CGF);
1390
1391 // Initialize the counter variable for the loop.
1392 QualType Int32Ty =
1393 CGF.getContext().getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/0);
1394 Address Counter = CGF.CreateMemTempWithoutCast(T: Int32Ty, Name: "critical_counter");
1395 LValue CounterLVal = CGF.MakeAddrLValue(Addr: Counter, T: Int32Ty);
1396 CGF.EmitStoreOfScalar(value: llvm::Constant::getNullValue(Ty: CGM.Int32Ty), lvalue: CounterLVal,
1397 /*isInit=*/true);
1398
1399 // Block checks if loop counter exceeds upper bound.
1400 CGF.EmitBlock(BB: LoopBB);
1401 llvm::Value *CounterVal = CGF.EmitLoadOfScalar(lvalue: CounterLVal, Loc);
1402 llvm::Value *CmpLoopBound = CGF.Builder.CreateICmpSLT(LHS: CounterVal, RHS: TeamWidth);
1403 CGF.Builder.CreateCondBr(Cond: CmpLoopBound, True: TestBB, False: ExitBB);
1404
1405 // Block tests which single thread should execute region, and which threads
1406 // should go straight to synchronisation point.
1407 CGF.EmitBlock(BB: TestBB);
1408 CounterVal = CGF.EmitLoadOfScalar(lvalue: CounterLVal, Loc);
1409 llvm::Value *CmpThreadToCounter =
1410 CGF.Builder.CreateICmpEQ(LHS: ThreadID, RHS: CounterVal);
1411 CGF.Builder.CreateCondBr(Cond: CmpThreadToCounter, True: BodyBB, False: SyncBB);
1412
1413 // Block emits the body of the critical region.
1414 CGF.EmitBlock(BB: BodyBB);
1415
1416 // Output the critical statement.
1417 CGOpenMPRuntime::emitCriticalRegion(CGF, CriticalName, CriticalOpGen, Loc,
1418 Hint);
1419
1420 // After the body surrounded by the critical region, the single executing
1421 // thread will jump to the synchronisation point.
1422 // Block waits for all threads in current team to finish then increments the
1423 // counter variable and returns to the loop.
1424 CGF.EmitBlock(BB: SyncBB);
1425 // Reconverge active threads in the warp.
1426 (void)CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1427 M&: CGM.getModule(), FnID: OMPRTL___kmpc_syncwarp),
1428 args: Mask);
1429
1430 llvm::Value *IncCounterVal =
1431 CGF.Builder.CreateNSWAdd(LHS: CounterVal, RHS: CGF.Builder.getInt32(C: 1));
1432 CGF.EmitStoreOfScalar(value: IncCounterVal, lvalue: CounterLVal);
1433 CGF.EmitBranch(Block: LoopBB);
1434
1435 // Block that is reached when all threads in the team complete the region.
1436 CGF.EmitBlock(BB: ExitBB, /*IsFinished=*/true);
1437}
1438
1439/// Cast value to the specified type.
1440static llvm::Value *castValueToType(CodeGenFunction &CGF, llvm::Value *Val,
1441 QualType ValTy, QualType CastTy,
1442 SourceLocation Loc) {
1443 assert(!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&
1444 "Cast type must sized.");
1445 assert(!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&
1446 "Val type must sized.");
1447 llvm::Type *LLVMCastTy = CGF.ConvertTypeForMem(T: CastTy);
1448 if (ValTy == CastTy)
1449 return Val;
1450 if (CGF.getContext().getTypeSizeInChars(T: ValTy) ==
1451 CGF.getContext().getTypeSizeInChars(T: CastTy))
1452 return CGF.Builder.CreateBitCast(V: Val, DestTy: LLVMCastTy);
1453 if (CastTy->isIntegerType() && ValTy->isIntegerType())
1454 return CGF.Builder.CreateIntCast(V: Val, DestTy: LLVMCastTy,
1455 isSigned: CastTy->hasSignedIntegerRepresentation());
1456 Address CastItem = CGF.CreateMemTempWithoutCast(T: CastTy);
1457 Address ValCastItem = CastItem.withElementType(ElemTy: Val->getType());
1458 CGF.EmitStoreOfScalar(Value: Val, Addr: ValCastItem, /*Volatile=*/false, Ty: ValTy,
1459 BaseInfo: LValueBaseInfo(AlignmentSource::Type),
1460 TBAAInfo: TBAAAccessInfo());
1461 return CGF.EmitLoadOfScalar(Addr: CastItem, /*Volatile=*/false, Ty: CastTy, Loc,
1462 BaseInfo: LValueBaseInfo(AlignmentSource::Type),
1463 TBAAInfo: TBAAAccessInfo());
1464}
1465
1466/// Extracts the built-in reduction operator from a combiner of the form `x = x
1467/// <op> rhs` (or the min/max conditional), or nullopt if the shape is not
1468/// recognized (e.g. user-defined reductions).
1469static std::optional<BinaryOperatorKind>
1470getReductionBinOpKind(const Expr *ReductionOp) {
1471 const auto *Assign = dyn_cast<BinaryOperator>(Val: ReductionOp);
1472 if (!Assign || Assign->getOpcode() != BO_Assign)
1473 return std::nullopt;
1474 const Expr *RHS = Assign->getRHS();
1475 // min/max are lowered as `x <cmp> rhs ? x : rhs`; the comparison identifies
1476 // it.
1477 if (const auto *ACO =
1478 dyn_cast<AbstractConditionalOperator>(Val: RHS->IgnoreParenImpCasts()))
1479 RHS = ACO->getCond();
1480 if (const auto *BO = dyn_cast<BinaryOperator>(Val: RHS->IgnoreParenImpCasts()))
1481 return BO->getOpcode();
1482 return std::nullopt;
1483}
1484
1485/// Maps a built-in reduction operator to an atomicrmw opcode for the atomic
1486/// cross-team reduction fast path, or nullopt if there is no direct atomicrmw
1487/// (e.g. user-defined, complex, fp min/max) so the buffer path is used instead.
1488static std::optional<llvm::AtomicRMWInst::BinOp>
1489getReductionAtomicRMWOp(BinaryOperatorKind BOK, QualType Ty) {
1490 bool IsInt = Ty->isIntegerType();
1491 bool IsSigned = Ty->hasSignedIntegerRepresentation();
1492 switch (BOK) {
1493 case BO_Add:
1494 case BO_Sub: // A `-` reduction sums the partials, so it accumulates with add.
1495 if (IsInt)
1496 return llvm::AtomicRMWInst::Add;
1497 if (Ty->isFloatingType())
1498 return llvm::AtomicRMWInst::FAdd;
1499 return std::nullopt;
1500 case BO_And:
1501 return IsInt ? std::optional(llvm::AtomicRMWInst::And) : std::nullopt;
1502 case BO_Or:
1503 return IsInt ? std::optional(llvm::AtomicRMWInst::Or) : std::nullopt;
1504 case BO_Xor:
1505 return IsInt ? std::optional(llvm::AtomicRMWInst::Xor) : std::nullopt;
1506 case BO_LT: // min
1507 if (IsInt)
1508 return IsSigned ? llvm::AtomicRMWInst::Min : llvm::AtomicRMWInst::UMin;
1509 return std::nullopt;
1510 case BO_GT: // max
1511 if (IsInt)
1512 return IsSigned ? llvm::AtomicRMWInst::Max : llvm::AtomicRMWInst::UMax;
1513 return std::nullopt;
1514 default:
1515 return std::nullopt;
1516 }
1517}
1518
1519///
1520/// Design of OpenMP reductions on the GPU
1521///
1522/// Consider a typical OpenMP program with one or more reduction
1523/// clauses:
1524///
1525/// float foo;
1526/// double bar;
1527/// #pragma omp target teams distribute parallel for \
1528/// reduction(+:foo) reduction(*:bar)
1529/// for (int i = 0; i < N; i++) {
1530/// foo += A[i]; bar *= B[i];
1531/// }
1532///
1533/// where 'foo' and 'bar' are reduced across all OpenMP threads in
1534/// all teams. In our OpenMP implementation on the NVPTX device an
1535/// OpenMP team is mapped to a CUDA threadblock and OpenMP threads
1536/// within a team are mapped to CUDA threads within a threadblock.
1537/// Our goal is to efficiently aggregate values across all OpenMP
1538/// threads such that:
1539///
1540/// - the compiler and runtime are logically concise, and
1541/// - the reduction is performed efficiently in a hierarchical
1542/// manner as follows: within OpenMP threads in the same warp,
1543/// across warps in a threadblock, and finally across teams on
1544/// the NVPTX device.
1545///
1546/// Introduction to Decoupling
1547///
1548/// We would like to decouple the compiler and the runtime so that the
1549/// latter is ignorant of the reduction variables (number, data types)
1550/// and the reduction operators. This allows a simpler interface
1551/// and implementation while still attaining good performance.
1552///
1553/// Pseudocode for the aforementioned OpenMP program generated by the
1554/// compiler is as follows:
1555///
1556/// 1. Create private copies of reduction variables on each OpenMP
1557/// thread: 'foo_private', 'bar_private'
1558/// 2. Each OpenMP thread reduces the chunk of 'A' and 'B' assigned
1559/// to it and writes the result in 'foo_private' and 'bar_private'
1560/// respectively.
1561/// 3. Call the OpenMP runtime on the GPU to reduce within a team
1562/// and store the result on the team master:
1563///
1564/// __kmpc_nvptx_parallel_reduce_nowait_v2(...,
1565/// reduceData, shuffleReduceFn, interWarpCpyFn)
1566///
1567/// where:
1568/// struct ReduceData {
1569/// double *foo;
1570/// double *bar;
1571/// } reduceData
1572/// reduceData.foo = &foo_private
1573/// reduceData.bar = &bar_private
1574///
1575/// 'shuffleReduceFn' and 'interWarpCpyFn' are pointers to two
1576/// auxiliary functions generated by the compiler that operate on
1577/// variables of type 'ReduceData'. They aid the runtime perform
1578/// algorithmic steps in a data agnostic manner.
1579///
1580/// 'shuffleReduceFn' is a pointer to a function that reduces data
1581/// of type 'ReduceData' across two OpenMP threads (lanes) in the
1582/// same warp. It takes the following arguments as input:
1583///
1584/// a. variable of type 'ReduceData' on the calling lane,
1585/// b. its lane_id,
1586/// c. an offset relative to the current lane_id to generate a
1587/// remote_lane_id. The remote lane contains the second
1588/// variable of type 'ReduceData' that is to be reduced.
1589/// d. an algorithm version parameter determining which reduction
1590/// algorithm to use.
1591///
1592/// 'shuffleReduceFn' retrieves data from the remote lane using
1593/// efficient GPU shuffle intrinsics and reduces, using the
1594/// algorithm specified by the 4th parameter, the two operands
1595/// element-wise. The result is written to the first operand.
1596///
1597/// Different reduction algorithms are implemented in different
1598/// runtime functions, all calling 'shuffleReduceFn' to perform
1599/// the essential reduction step. Therefore, based on the 4th
1600/// parameter, this function behaves slightly differently to
1601/// cooperate with the runtime to ensure correctness under
1602/// different circumstances.
1603///
1604/// 'InterWarpCpyFn' is a pointer to a function that transfers
1605/// reduced variables across warps. It tunnels, through CUDA
1606/// shared memory, the thread-private data of type 'ReduceData'
1607/// from lane 0 of each warp to a lane in the first warp.
1608/// 4. Call the OpenMP runtime on the GPU to reduce across teams.
1609/// The last team writes the global reduced value to memory.
1610///
1611/// ret = __kmpc_nvptx_teams_reduce_nowait(...,
1612/// reduceData, shuffleReduceFn, interWarpCpyFn,
1613/// scratchpadCopyFn, loadAndReduceFn)
1614///
1615/// 'scratchpadCopyFn' is a helper that stores reduced
1616/// data from the team master to a scratchpad array in
1617/// global memory.
1618///
1619/// 'loadAndReduceFn' is a helper that loads data from
1620/// the scratchpad array and reduces it with the input
1621/// operand.
1622///
1623/// These compiler generated functions hide address
1624/// calculation and alignment information from the runtime.
1625/// 5. if ret == 1:
1626/// The team master of the last team stores the reduced
1627/// result to the globals in memory.
1628/// foo += reduceData.foo; bar *= reduceData.bar
1629///
1630///
1631/// Warp Reduction Algorithms
1632///
1633/// On the warp level, we have three algorithms implemented in the
1634/// OpenMP runtime depending on the number of active lanes:
1635///
1636/// Full Warp Reduction
1637///
1638/// The reduce algorithm within a warp where all lanes are active
1639/// is implemented in the runtime as follows:
1640///
1641/// full_warp_reduce(void *reduce_data,
1642/// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
1643/// for (int offset = WARPSIZE/2; offset > 0; offset /= 2)
1644/// ShuffleReduceFn(reduce_data, 0, offset, 0);
1645/// }
1646///
1647/// The algorithm completes in log(2, WARPSIZE) steps.
1648///
1649/// 'ShuffleReduceFn' is used here with lane_id set to 0 because it is
1650/// not used therefore we save instructions by not retrieving lane_id
1651/// from the corresponding special registers. The 4th parameter, which
1652/// represents the version of the algorithm being used, is set to 0 to
1653/// signify full warp reduction.
1654///
1655/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
1656///
1657/// #reduce_elem refers to an element in the local lane's data structure
1658/// #remote_elem is retrieved from a remote lane
1659/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
1660/// reduce_elem = reduce_elem REDUCE_OP remote_elem;
1661///
1662/// Contiguous Partial Warp Reduction
1663///
1664/// This reduce algorithm is used within a warp where only the first
1665/// 'n' (n <= WARPSIZE) lanes are active. It is typically used when the
1666/// number of OpenMP threads in a parallel region is not a multiple of
1667/// WARPSIZE. The algorithm is implemented in the runtime as follows:
1668///
1669/// void
1670/// contiguous_partial_reduce(void *reduce_data,
1671/// kmp_ShuffleReductFctPtr ShuffleReduceFn,
1672/// int size, int lane_id) {
1673/// int curr_size;
1674/// int offset;
1675/// curr_size = size;
1676/// mask = curr_size/2;
1677/// while (offset>0) {
1678/// ShuffleReduceFn(reduce_data, lane_id, offset, 1);
1679/// curr_size = (curr_size+1)/2;
1680/// offset = curr_size/2;
1681/// }
1682/// }
1683///
1684/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
1685///
1686/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
1687/// if (lane_id < offset)
1688/// reduce_elem = reduce_elem REDUCE_OP remote_elem
1689/// else
1690/// reduce_elem = remote_elem
1691///
1692/// This algorithm assumes that the data to be reduced are located in a
1693/// contiguous subset of lanes starting from the first. When there is
1694/// an odd number of active lanes, the data in the last lane is not
1695/// aggregated with any other lane's dat but is instead copied over.
1696///
1697/// Dispersed Partial Warp Reduction
1698///
1699/// This algorithm is used within a warp when any discontiguous subset of
1700/// lanes are active. It is used to implement the reduction operation
1701/// across lanes in an OpenMP simd region or in a nested parallel region.
1702///
1703/// void
1704/// dispersed_partial_reduce(void *reduce_data,
1705/// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
1706/// int size, remote_id;
1707/// int logical_lane_id = number_of_active_lanes_before_me() * 2;
1708/// do {
1709/// remote_id = next_active_lane_id_right_after_me();
1710/// # the above function returns 0 of no active lane
1711/// # is present right after the current lane.
1712/// size = number_of_active_lanes_in_this_warp();
1713/// logical_lane_id /= 2;
1714/// ShuffleReduceFn(reduce_data, logical_lane_id,
1715/// remote_id-1-threadIdx.x, 2);
1716/// } while (logical_lane_id % 2 == 0 && size > 1);
1717/// }
1718///
1719/// There is no assumption made about the initial state of the reduction.
1720/// Any number of lanes (>=1) could be active at any position. The reduction
1721/// result is returned in the first active lane.
1722///
1723/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
1724///
1725/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
1726/// if (lane_id % 2 == 0 && offset > 0)
1727/// reduce_elem = reduce_elem REDUCE_OP remote_elem
1728/// else
1729/// reduce_elem = remote_elem
1730///
1731///
1732/// Intra-Team Reduction
1733///
1734/// This function, as implemented in the runtime call
1735/// '__kmpc_nvptx_parallel_reduce_nowait_v2', aggregates data across OpenMP
1736/// threads in a team. It first reduces within a warp using the
1737/// aforementioned algorithms. We then proceed to gather all such
1738/// reduced values at the first warp.
1739///
1740/// The runtime makes use of the function 'InterWarpCpyFn', which copies
1741/// data from each of the "warp master" (zeroth lane of each warp, where
1742/// warp-reduced data is held) to the zeroth warp. This step reduces (in
1743/// a mathematical sense) the problem of reduction across warp masters in
1744/// a block to the problem of warp reduction.
1745///
1746///
1747/// Inter-Team Reduction
1748///
1749/// Once a team has reduced its data to a single value, it is stored in
1750/// a global scratchpad array. Since each team has a distinct slot, this
1751/// can be done without locking.
1752///
1753/// The last team to write to the scratchpad array proceeds to reduce the
1754/// scratchpad array. One or more workers in the last team use the helper
1755/// 'loadAndReduceDataFn' to load and reduce values from the array, i.e.,
1756/// the k'th worker reduces every k'th element.
1757///
1758/// Finally, a call is made to '__kmpc_nvptx_parallel_reduce_nowait_v2' to
1759/// reduce across workers and compute a globally reduced value.
1760///
1761void CGOpenMPRuntimeGPU::emitReduction(
1762 CodeGenFunction &CGF, SourceLocation Loc, ArrayRef<const Expr *> Privates,
1763 ArrayRef<const Expr *> LHSExprs, ArrayRef<const Expr *> RHSExprs,
1764 ArrayRef<const Expr *> ReductionOps, ReductionOptionsTy Options) {
1765 if (!CGF.HaveInsertPoint())
1766 return;
1767
1768 bool ParallelReduction = isOpenMPParallelDirective(DKind: Options.ReductionKind);
1769 bool TeamsReduction = isOpenMPTeamsDirective(DKind: Options.ReductionKind);
1770
1771 if (Options.SimpleReduction) {
1772 assert(!TeamsReduction && !ParallelReduction &&
1773 "Invalid reduction selection in emitReduction.");
1774 (void)ParallelReduction;
1775 CGOpenMPRuntime::emitReduction(CGF, Loc, Privates, LHSExprs, RHSExprs,
1776 ReductionOps, Options);
1777 return;
1778 }
1779
1780 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> VarFieldMap;
1781 llvm::SmallVector<const ValueDecl *, 4> PrivatesReductions(Privates.size());
1782 int Cnt = 0;
1783 for (const Expr *DRE : Privates) {
1784 PrivatesReductions[Cnt] = cast<DeclRefExpr>(Val: DRE)->getDecl();
1785 ++Cnt;
1786 }
1787 const RecordDecl *ReductionRec = ::buildRecordForGlobalizedVars(
1788 C&: CGM.getContext(), EscapedDecls: PrivatesReductions, EscapedDeclsForTeams: {}, MappedDeclsFields&: VarFieldMap, BufSize: 1);
1789
1790 // The atomic cross-team reduction fast path is opt-in. Hand each eligible
1791 // scalar reduction an atomic combiner; createReductionsGPU uses the atomic
1792 // path only if every reduction in the set has one. Track whether that holds
1793 // so we can skip the (then unused) per-team buffer registration.
1794 bool UseAtomicReduction =
1795 TeamsReduction && CGM.getLangOpts().OpenMPTargetAtomicReduction;
1796 bool AllAtomicable = UseAtomicReduction;
1797
1798 // Source location for the ident struct
1799 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
1800
1801 using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
1802 InsertPointTy AllocaIP(CGF.AllocaInsertPt->getIterator());
1803 InsertPointTy CodeGenIP(CGF.Builder.GetInsertPoint());
1804 llvm::OpenMPIRBuilder::LocationDescription OmpLoc(
1805 CodeGenIP, CGF.SourceLocToDebugLoc(Location: Loc));
1806 llvm::SmallVector<llvm::OpenMPIRBuilder::ReductionInfo, 2> ReductionInfos;
1807
1808 CodeGenFunction::OMPPrivateScope Scope(CGF);
1809 unsigned Idx = 0;
1810 for (const Expr *Private : Privates) {
1811 llvm::Type *ElementType;
1812 llvm::Value *Variable;
1813 llvm::Value *PrivateVariable;
1814 llvm::OpenMPIRBuilder::ReductionGenAtomicCBTy AtomicReductionGen = nullptr;
1815 ElementType = CGF.ConvertTypeForMem(T: Private->getType());
1816 const auto *RHSVar =
1817 cast<VarDecl>(Val: cast<DeclRefExpr>(Val: RHSExprs[Idx])->getDecl());
1818 PrivateVariable = CGF.GetAddrOfLocalVar(VD: RHSVar).emitRawPointer(CGF);
1819 const auto *LHSVar =
1820 cast<VarDecl>(Val: cast<DeclRefExpr>(Val: LHSExprs[Idx])->getDecl());
1821 Variable = CGF.GetAddrOfLocalVar(VD: LHSVar).emitRawPointer(CGF);
1822 llvm::OpenMPIRBuilder::EvalKind EvalKind;
1823 switch (CGF.getEvaluationKind(T: Private->getType())) {
1824 case TEK_Scalar:
1825 EvalKind = llvm::OpenMPIRBuilder::EvalKind::Scalar;
1826 break;
1827 case TEK_Complex:
1828 EvalKind = llvm::OpenMPIRBuilder::EvalKind::Complex;
1829 break;
1830 case TEK_Aggregate:
1831 EvalKind = llvm::OpenMPIRBuilder::EvalKind::Aggregate;
1832 break;
1833 }
1834 auto ReductionGen = [&](InsertPointTy CodeGenIP, unsigned I,
1835 llvm::Value **LHSPtr, llvm::Value **RHSPtr,
1836 llvm::Function *NewFunc) {
1837 CGF.Builder.restoreIP(IP: CodeGenIP);
1838 auto *CurFn = CGF.CurFn;
1839 CGF.CurFn = NewFunc;
1840
1841 // The helper has no DISubprogram of its own, so a debug location here
1842 // would name the enclosing function's scope, which is invalid IR.
1843 // Suppress them, as the other OpenMPIRBuilder-generated helpers do.
1844 llvm::DebugLoc SavedDebugLoc = CGF.Builder.getCurrentDebugLocation();
1845 CGF.Builder.SetCurrentDebugLocation(llvm::DebugLoc());
1846 CGF.disableDebugInfo();
1847
1848 *LHSPtr = CGF.GetAddrOfLocalVar(
1849 VD: cast<VarDecl>(Val: cast<DeclRefExpr>(Val: LHSExprs[I])->getDecl()))
1850 .emitRawPointer(CGF);
1851 *RHSPtr = CGF.GetAddrOfLocalVar(
1852 VD: cast<VarDecl>(Val: cast<DeclRefExpr>(Val: RHSExprs[I])->getDecl()))
1853 .emitRawPointer(CGF);
1854
1855 emitSingleReductionCombiner(CGF, ReductionOp: ReductionOps[I], PrivateRef: Privates[I],
1856 LHS: cast<DeclRefExpr>(Val: LHSExprs[I]),
1857 RHS: cast<DeclRefExpr>(Val: RHSExprs[I]));
1858
1859 CGF.enableDebugInfo();
1860 CGF.Builder.SetCurrentDebugLocation(SavedDebugLoc);
1861 CGF.CurFn = CurFn;
1862
1863 return CGF.Builder.GetInsertPoint();
1864 };
1865
1866 // For the atomic fast path, hand this reduction an atomic combiner if it is
1867 // a scalar with a direct atomicrmw; otherwise the set is not fully
1868 // atomicable and falls back to the buffer path.
1869 if (UseAtomicReduction) {
1870 std::optional<llvm::AtomicRMWInst::BinOp> AtomicOp;
1871 if (EvalKind == llvm::OpenMPIRBuilder::EvalKind::Scalar) {
1872 if (std::optional<BinaryOperatorKind> BOK =
1873 getReductionBinOpKind(ReductionOp: ReductionOps[Idx]))
1874 AtomicOp = getReductionAtomicRMWOp(BOK: *BOK, Ty: Private->getType());
1875 }
1876 if (!AtomicOp) {
1877 AllAtomicable = false;
1878 } else {
1879 llvm::AtomicRMWInst::BinOp Op = *AtomicOp;
1880 llvm::Align Alignment =
1881 CGM.getModule().getDataLayout().getPrefTypeAlign(Ty: ElementType);
1882 // Device (agent) scope suffices: all teams accumulate on-device and the
1883 // host reads the result only after the kernel (via map-back), so the
1884 // far costlier system scope is unnecessary. The
1885 // no.fine.grained/no.remote memory metadata is omitted so the atomic
1886 // stays correct under USM.
1887 llvm::SyncScope::ID SSID = CGF.getTargetHooks().getLLVMSyncScopeID(
1888 LangOpts: CGF.getLangOpts(), Scope: SyncScope::DeviceScope,
1889 Ordering: llvm::AtomicOrdering::Monotonic, Ctx&: CGF.getLLVMContext());
1890 AtomicReductionGen = [Op, Alignment,
1891 SSID](InsertPointTy IP, llvm::Type *EltTy,
1892 llvm::Value *LHS, llvm::Value *RHS)
1893 -> llvm::OpenMPIRBuilder::InsertPointOrErrorTy {
1894 llvm::IRBuilder<> Builder(IP);
1895 llvm::Value *Val = Builder.CreateLoad(Ty: EltTy, Ptr: RHS);
1896 Builder.CreateAtomicRMW(Op, Ptr: LHS, Val, Align: Alignment,
1897 Ordering: llvm::AtomicOrdering::Monotonic, SSID);
1898 return Builder.GetInsertPoint();
1899 };
1900 }
1901 }
1902
1903 ReductionInfos.emplace_back(Args: llvm::OpenMPIRBuilder::ReductionInfo(
1904 ElementType, Variable, PrivateVariable, EvalKind,
1905 /*ReductionGen=*/nullptr, ReductionGen, AtomicReductionGen,
1906 /*DataPtrPtrGen=*/nullptr));
1907 Idx++;
1908 }
1909
1910 // The atomic path folds directly into the mapped variable and needs no
1911 // per-team buffer; register the record for buffer allocation otherwise.
1912 if (TeamsReduction && !AllAtomicable)
1913 TeamsReductions.push_back(Elt: ReductionRec);
1914
1915 bool IsSPMD = getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD;
1916 llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
1917 cantFail(ValOrErr: OMPBuilder.createReductionsGPU(
1918 Loc: OmpLoc, AllocaIP, CodeGenIP, ReductionInfos, /*IsByRef=*/{}, IsNoWait: false,
1919 IsTeamsReduction: TeamsReduction, IsSPMD,
1920 ReductionGenCBKind: llvm::OpenMPIRBuilder::ReductionGenCBKind::Clang,
1921 GridValue: CGF.getTarget().getGridValue(), SrcLocInfo: RTLoc));
1922 CGF.Builder.restoreIP(IP: AfterIP);
1923}
1924
1925const VarDecl *
1926CGOpenMPRuntimeGPU::translateParameter(const FieldDecl *FD,
1927 const VarDecl *NativeParam) const {
1928 if (!NativeParam->getType()->isReferenceType())
1929 return NativeParam;
1930 QualType ArgType = NativeParam->getType();
1931 QualifierCollector QC;
1932 const Type *NonQualTy = QC.strip(type: ArgType);
1933 QualType PointeeTy = cast<ReferenceType>(Val: NonQualTy)->getPointeeType();
1934 if (const auto *Attr = FD->getAttr<OMPCaptureKindAttr>()) {
1935 if (Attr->getCaptureKind() == OMPC_map) {
1936 PointeeTy = CGM.getContext().getAddrSpaceQualType(T: PointeeTy,
1937 AddressSpace: LangAS::opencl_global);
1938 }
1939 }
1940 ArgType = CGM.getContext().getPointerType(T: PointeeTy);
1941 QC.addRestrict();
1942 ArgType = QC.apply(Context: CGM.getContext(), QT: ArgType);
1943 if (isa<ImplicitParamDecl>(Val: NativeParam))
1944 return ImplicitParamDecl::Create(
1945 C&: CGM.getContext(), /*DC=*/nullptr, IdLoc: NativeParam->getLocation(),
1946 Id: NativeParam->getIdentifier(), T: ArgType, ParamKind: ImplicitParamKind::Other);
1947 return ParmVarDecl::Create(
1948 C&: CGM.getContext(),
1949 DC: const_cast<DeclContext *>(NativeParam->getDeclContext()),
1950 StartLoc: NativeParam->getBeginLoc(), IdLoc: NativeParam->getLocation(),
1951 Id: NativeParam->getIdentifier(), T: ArgType,
1952 /*TInfo=*/nullptr, S: SC_None, /*DefArg=*/nullptr);
1953}
1954
1955Address
1956CGOpenMPRuntimeGPU::getParameterAddress(CodeGenFunction &CGF,
1957 const VarDecl *NativeParam,
1958 const VarDecl *TargetParam) const {
1959 assert(NativeParam != TargetParam &&
1960 NativeParam->getType()->isReferenceType() &&
1961 "Native arg must not be the same as target arg.");
1962 Address LocalAddr = CGF.GetAddrOfLocalVar(VD: TargetParam);
1963 QualType NativeParamType = NativeParam->getType();
1964 QualifierCollector QC;
1965 const Type *NonQualTy = QC.strip(type: NativeParamType);
1966 QualType NativePointeeTy = cast<ReferenceType>(Val: NonQualTy)->getPointeeType();
1967 unsigned NativePointeeAddrSpace =
1968 CGF.getTypes().getTargetAddressSpace(T: NativePointeeTy);
1969 QualType TargetTy = TargetParam->getType();
1970 llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(Addr: LocalAddr, /*Volatile=*/false,
1971 Ty: TargetTy, Loc: SourceLocation());
1972 // Cast to native address space.
1973 TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
1974 V: TargetAddr,
1975 DestTy: llvm::PointerType::get(C&: CGF.getLLVMContext(), AddressSpace: NativePointeeAddrSpace));
1976 Address NativeParamAddr = CGF.CreateMemTemp(T: NativeParamType);
1977 CGF.EmitStoreOfScalar(Value: TargetAddr, Addr: NativeParamAddr, /*Volatile=*/false,
1978 Ty: NativeParamType);
1979 return NativeParamAddr;
1980}
1981
1982void CGOpenMPRuntimeGPU::emitOutlinedFunctionCall(
1983 CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn,
1984 ArrayRef<llvm::Value *> Args) const {
1985 SmallVector<llvm::Value *, 4> TargetArgs;
1986 TargetArgs.reserve(N: Args.size());
1987 auto *FnType = OutlinedFn.getFunctionType();
1988 for (unsigned I = 0, E = Args.size(); I < E; ++I) {
1989 if (FnType->isVarArg() && FnType->getNumParams() <= I) {
1990 TargetArgs.append(in_start: std::next(x: Args.begin(), n: I), in_end: Args.end());
1991 break;
1992 }
1993 llvm::Type *TargetType = FnType->getParamType(i: I);
1994 llvm::Value *NativeArg = Args[I];
1995 if (!TargetType->isPointerTy()) {
1996 TargetArgs.emplace_back(Args&: NativeArg);
1997 continue;
1998 }
1999 TargetArgs.emplace_back(
2000 Args: CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(V: NativeArg, DestTy: TargetType));
2001 }
2002 CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, Args: TargetArgs);
2003}
2004
2005/// Emit function which wraps the outline parallel region
2006/// and controls the arguments which are passed to this function.
2007/// The wrapper ensures that the outlined function is called
2008/// with the correct arguments when data is shared.
2009llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
2010 llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) {
2011 ASTContext &Ctx = CGM.getContext();
2012 const auto &CS = *D.getCapturedStmt(RegionKind: OMPD_parallel);
2013
2014 // Create a function that takes as argument the source thread.
2015 FunctionArgList WrapperArgs;
2016 QualType Int16QTy =
2017 Ctx.getIntTypeForBitwidth(/*DestWidth=*/16, /*Signed=*/false);
2018 QualType Int32QTy =
2019 Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/false);
2020 auto *ParallelLevelArg = ImplicitParamDecl::Create(
2021 C&: Ctx, /*DC=*/nullptr, IdLoc: D.getBeginLoc(),
2022 /*Id=*/nullptr, T: Int16QTy, ParamKind: ImplicitParamKind::Other);
2023 auto *WrapperArg = ImplicitParamDecl::Create(
2024 C&: Ctx, /*DC=*/nullptr, IdLoc: D.getBeginLoc(),
2025 /*Id=*/nullptr, T: Int32QTy, ParamKind: ImplicitParamKind::Other);
2026 WrapperArgs.emplace_back(Args&: ParallelLevelArg);
2027 WrapperArgs.emplace_back(Args&: WrapperArg);
2028
2029 const CGFunctionInfo &CGFI =
2030 CGM.getTypes().arrangeBuiltinFunctionDeclaration(resultType: Ctx.VoidTy, args: WrapperArgs);
2031
2032 auto *Fn = llvm::Function::Create(
2033 Ty: CGM.getTypes().GetFunctionType(Info: CGFI), Linkage: llvm::GlobalValue::InternalLinkage,
2034 N: Twine(OutlinedParallelFn->getName(), "_wrapper"), M: &CGM.getModule());
2035
2036 // Ensure we do not inline the function. This is trivially true for the ones
2037 // passed to __kmpc_fork_call but the ones calles in serialized regions
2038 // could be inlined. This is not a perfect but it is closer to the invariant
2039 // we want, namely, every data environment starts with a new function.
2040 // TODO: We should pass the if condition to the runtime function and do the
2041 // handling there. Much cleaner code.
2042 Fn->addFnAttr(Kind: llvm::Attribute::NoInline);
2043
2044 CGM.SetInternalFunctionAttributes(GD: GlobalDecl(), F: Fn, FI: CGFI);
2045 Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
2046
2047 CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
2048 CGF.StartFunction(GD: GlobalDecl(), RetTy: Ctx.VoidTy, Fn, FnInfo: CGFI, Args: WrapperArgs,
2049 Loc: D.getBeginLoc(), StartLoc: D.getBeginLoc());
2050
2051 const auto *RD = CS.getCapturedRecordDecl();
2052 auto CurField = RD->field_begin();
2053
2054 Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(Ty: CGF.Int32Ty,
2055 /*Name=*/".zero.addr");
2056 CGF.Builder.CreateStore(Val: CGF.Builder.getInt32(/*C*/ 0), Addr: ZeroAddr);
2057 // Get the array of arguments.
2058 SmallVector<llvm::Value *, 8> Args;
2059
2060 Args.emplace_back(Args: CGF.GetAddrOfLocalVar(VD: WrapperArg).emitRawPointer(CGF));
2061 Args.emplace_back(Args: ZeroAddr.emitRawPointer(CGF));
2062
2063 CGBuilderTy &Bld = CGF.Builder;
2064 auto CI = CS.capture_begin();
2065
2066 // Use global memory for data sharing.
2067 // Handle passing of global args to workers.
2068 RawAddress GlobalArgs =
2069 CGF.CreateDefaultAlignTempAlloca(Ty: CGF.VoidPtrPtrTy, Name: "global_args");
2070 llvm::Value *GlobalArgsPtr = GlobalArgs.getPointer();
2071 llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};
2072 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
2073 M&: CGM.getModule(), FnID: OMPRTL___kmpc_get_shared_variables),
2074 args: DataSharingArgs);
2075
2076 // Retrieve the shared variables from the list of references returned
2077 // by the runtime. Pass the variables to the outlined function.
2078 Address SharedArgListAddress = Address::invalid();
2079 if (CS.capture_size() > 0 ||
2080 isOpenMPLoopBoundSharingDirective(Kind: D.getDirectiveKind())) {
2081 SharedArgListAddress = CGF.EmitLoadOfPointer(
2082 Ptr: GlobalArgs, PtrTy: CGF.getContext()
2083 .getPointerType(T: CGF.getContext().VoidPtrTy)
2084 .castAs<PointerType>());
2085 }
2086 unsigned Idx = 0;
2087 if (isOpenMPLoopBoundSharingDirective(Kind: D.getDirectiveKind())) {
2088 Address Src = Bld.CreateConstInBoundsGEP(Addr: SharedArgListAddress, Index: Idx);
2089 Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
2090 Addr: Src, Ty: Bld.getPtrTy(AddrSpace: 0), ElementTy: CGF.SizeTy);
2091 llvm::Value *LB = CGF.EmitLoadOfScalar(
2092 Addr: TypedAddress,
2093 /*Volatile=*/false,
2094 Ty: CGF.getContext().getPointerType(T: CGF.getContext().getSizeType()),
2095 Loc: cast<OMPLoopDirective>(Val: D).getLowerBoundVariable()->getExprLoc());
2096 Args.emplace_back(Args&: LB);
2097 ++Idx;
2098 Src = Bld.CreateConstInBoundsGEP(Addr: SharedArgListAddress, Index: Idx);
2099 TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(Addr: Src, Ty: Bld.getPtrTy(AddrSpace: 0),
2100 ElementTy: CGF.SizeTy);
2101 llvm::Value *UB = CGF.EmitLoadOfScalar(
2102 Addr: TypedAddress,
2103 /*Volatile=*/false,
2104 Ty: CGF.getContext().getPointerType(T: CGF.getContext().getSizeType()),
2105 Loc: cast<OMPLoopDirective>(Val: D).getUpperBoundVariable()->getExprLoc());
2106 Args.emplace_back(Args&: UB);
2107 ++Idx;
2108 }
2109 if (CS.capture_size() > 0) {
2110 ASTContext &CGFContext = CGF.getContext();
2111 for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CI, ++CurField) {
2112 QualType ElemTy = CurField->getType();
2113 Address Src = Bld.CreateConstInBoundsGEP(Addr: SharedArgListAddress, Index: I + Idx);
2114 Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
2115 Addr: Src, Ty: CGF.ConvertTypeForMem(T: CGFContext.getPointerType(T: ElemTy)),
2116 ElementTy: CGF.ConvertTypeForMem(T: ElemTy));
2117 llvm::Value *Arg = CGF.EmitLoadOfScalar(Addr: TypedAddress,
2118 /*Volatile=*/false,
2119 Ty: CGFContext.getPointerType(T: ElemTy),
2120 Loc: CI->getLocation());
2121 if (CI->capturesVariableByCopy() &&
2122 !CI->getCapturedVar()->getType()->isAnyPointerType()) {
2123 Arg = castValueToType(CGF, Val: Arg, ValTy: ElemTy, CastTy: CGFContext.getUIntPtrType(),
2124 Loc: CI->getLocation());
2125 }
2126 Args.emplace_back(Args&: Arg);
2127 }
2128 }
2129
2130 emitOutlinedFunctionCall(CGF, Loc: D.getBeginLoc(), OutlinedFn: OutlinedParallelFn, Args);
2131 CGF.FinishFunction();
2132 return Fn;
2133}
2134
2135void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,
2136 const Decl *D) {
2137 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
2138 return;
2139
2140 assert(D && "Expected function or captured|block decl.");
2141 assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&
2142 "Function is registered already.");
2143 assert((!TeamAndReductions.first || TeamAndReductions.first == D) &&
2144 "Team is set but not processed.");
2145 const Stmt *Body = nullptr;
2146 bool NeedToDelayGlobalization = false;
2147 if (const auto *FD = dyn_cast<FunctionDecl>(Val: D)) {
2148 Body = FD->getBody();
2149 } else if (const auto *BD = dyn_cast<BlockDecl>(Val: D)) {
2150 Body = BD->getBody();
2151 } else if (const auto *CD = dyn_cast<CapturedDecl>(Val: D)) {
2152 Body = CD->getBody();
2153 NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;
2154 if (NeedToDelayGlobalization &&
2155 getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
2156 return;
2157 }
2158 if (!Body)
2159 return;
2160 CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);
2161 VarChecker.Visit(S: Body);
2162 const RecordDecl *GlobalizedVarsRecord =
2163 VarChecker.getGlobalizedRecord(IsInTTDRegion);
2164 TeamAndReductions.first = nullptr;
2165 TeamAndReductions.second.clear();
2166 ArrayRef<const ValueDecl *> EscapedVariableLengthDecls =
2167 VarChecker.getEscapedVariableLengthDecls();
2168 ArrayRef<const ValueDecl *> DelayedVariableLengthDecls =
2169 VarChecker.getDelayedVariableLengthDecls();
2170 if (!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty() &&
2171 DelayedVariableLengthDecls.empty())
2172 return;
2173 auto I = FunctionGlobalizedDecls.try_emplace(Key: CGF.CurFn).first;
2174 I->getSecond().MappedParams =
2175 std::make_unique<CodeGenFunction::OMPMapVars>();
2176 I->getSecond().EscapedParameters.insert(
2177 I: VarChecker.getEscapedParameters().begin(),
2178 E: VarChecker.getEscapedParameters().end());
2179 I->getSecond().EscapedVariableLengthDecls.append(
2180 in_start: EscapedVariableLengthDecls.begin(), in_end: EscapedVariableLengthDecls.end());
2181 I->getSecond().DelayedVariableLengthDecls.append(
2182 in_start: DelayedVariableLengthDecls.begin(), in_end: DelayedVariableLengthDecls.end());
2183 DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
2184 for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
2185 assert(VD->isCanonicalDecl() && "Expected canonical declaration");
2186 Data.try_emplace(Key: VD);
2187 }
2188 if (!NeedToDelayGlobalization) {
2189 emitGenericVarsProlog(CGF, Loc: D->getBeginLoc());
2190 struct GlobalizationScope final : EHScopeStack::Cleanup {
2191 GlobalizationScope() = default;
2192
2193 void Emit(CodeGenFunction &CGF, Flags flags) override {
2194 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
2195 .emitGenericVarsEpilog(CGF);
2196 }
2197 };
2198 CGF.EHStack.pushCleanup<GlobalizationScope>(Kind: NormalAndEHCleanup);
2199 }
2200}
2201
2202Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
2203 const VarDecl *VD) {
2204 if (VD && VD->hasAttr<OMPAllocateDeclAttr>()) {
2205 const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
2206 auto AS = LangAS::Default;
2207 switch (A->getAllocatorType()) {
2208 case OMPAllocateDeclAttr::OMPNullMemAlloc:
2209 case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
2210 case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
2211 case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
2212 break;
2213 case OMPAllocateDeclAttr::OMPThreadMemAlloc:
2214 return Address::invalid();
2215 case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
2216 // TODO: implement aupport for user-defined allocators.
2217 return Address::invalid();
2218 case OMPAllocateDeclAttr::OMPConstMemAlloc:
2219 AS = LangAS::cuda_constant;
2220 break;
2221 case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
2222 AS = LangAS::cuda_shared;
2223 break;
2224 case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
2225 case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
2226 break;
2227 }
2228 llvm::Type *VarTy = CGF.ConvertTypeForMem(T: VD->getType());
2229 auto *GV = new llvm::GlobalVariable(
2230 CGM.getModule(), VarTy, /*isConstant=*/false,
2231 llvm::GlobalValue::InternalLinkage, llvm::PoisonValue::get(T: VarTy),
2232 VD->getName(),
2233 /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
2234 CGM.getContext().getTargetAddressSpace(AS));
2235 CharUnits Align = CGM.getContext().getDeclAlign(D: VD);
2236 GV->setAlignment(Align.getAsAlign());
2237 return Address(
2238 CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
2239 V: GV, DestTy: CGF.Builder.getPtrTy(AddrSpace: CGM.getContext().getTargetAddressSpace(
2240 AS: VD->getType().getAddressSpace()))),
2241 VarTy, Align);
2242 }
2243
2244 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
2245 return Address::invalid();
2246
2247 VD = VD->getCanonicalDecl();
2248 auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
2249 if (I == FunctionGlobalizedDecls.end())
2250 return Address::invalid();
2251 auto VDI = I->getSecond().LocalVarData.find(Key: VD);
2252 if (VDI != I->getSecond().LocalVarData.end())
2253 return VDI->second.PrivateAddr;
2254 if (VD->hasAttrs()) {
2255 for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),
2256 E(VD->attr_end());
2257 IT != E; ++IT) {
2258 auto VDI = I->getSecond().LocalVarData.find(
2259 Key: cast<VarDecl>(Val: cast<DeclRefExpr>(Val: IT->getRef())->getDecl())
2260 ->getCanonicalDecl());
2261 if (VDI != I->getSecond().LocalVarData.end())
2262 return VDI->second.PrivateAddr;
2263 }
2264 }
2265
2266 return Address::invalid();
2267}
2268
2269void CGOpenMPRuntimeGPU::functionFinished(CodeGenFunction &CGF) {
2270 FunctionGlobalizedDecls.erase(Val: CGF.CurFn);
2271 CGOpenMPRuntime::functionFinished(CGF);
2272}
2273
2274void CGOpenMPRuntimeGPU::getDefaultDistScheduleAndChunk(
2275 CodeGenFunction &CGF, const OMPLoopDirective &S,
2276 OpenMPDistScheduleClauseKind &ScheduleKind,
2277 llvm::Value *&Chunk) const {
2278 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
2279 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
2280 ScheduleKind = OMPC_DIST_SCHEDULE_static;
2281 Chunk = CGF.EmitScalarConversion(
2282 Src: RT.getGPUNumThreads(CGF),
2283 SrcTy: CGF.getContext().getIntTypeForBitwidth(DestWidth: 32, /*Signed=*/0),
2284 DstTy: S.getIterationVariable()->getType(), Loc: S.getBeginLoc());
2285 return;
2286 }
2287 CGOpenMPRuntime::getDefaultDistScheduleAndChunk(
2288 CGF, S, ScheduleKind, Chunk);
2289}
2290
2291void CGOpenMPRuntimeGPU::getDefaultScheduleAndChunk(
2292 CodeGenFunction &CGF, const OMPLoopDirective &S,
2293 OpenMPScheduleClauseKind &ScheduleKind,
2294 const Expr *&ChunkExpr) const {
2295 ScheduleKind = OMPC_SCHEDULE_static;
2296 // Chunk size is 1 in this case.
2297 llvm::APInt ChunkSize(32, 1);
2298 ChunkExpr = IntegerLiteral::Create(C: CGF.getContext(), V: ChunkSize,
2299 type: CGF.getContext().getIntTypeForBitwidth(DestWidth: 32, /*Signed=*/0),
2300 l: SourceLocation());
2301}
2302
2303void CGOpenMPRuntimeGPU::adjustTargetSpecificDataForLambdas(
2304 CodeGenFunction &CGF, const OMPExecutableDirective &D) const {
2305 assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
2306 " Expected target-based directive.");
2307 const CapturedStmt *CS = D.getCapturedStmt(RegionKind: OMPD_target);
2308 for (const CapturedStmt::Capture &C : CS->captures()) {
2309 // Capture variables captured by reference in lambdas for target-based
2310 // directives.
2311 if (!C.capturesVariable())
2312 continue;
2313 const VarDecl *VD = C.getCapturedVar();
2314 const auto *RD = VD->getType()
2315 .getCanonicalType()
2316 .getNonReferenceType()
2317 ->getAsCXXRecordDecl();
2318 if (!RD || !RD->isLambda())
2319 continue;
2320 Address VDAddr = CGF.GetAddrOfLocalVar(VD);
2321 LValue VDLVal;
2322 if (VD->getType().getCanonicalType()->isReferenceType())
2323 VDLVal = CGF.EmitLoadOfReferenceLValue(RefAddr: VDAddr, RefTy: VD->getType());
2324 else
2325 VDLVal = CGF.MakeAddrLValue(
2326 Addr: VDAddr, T: VD->getType().getCanonicalType().getNonReferenceType());
2327 llvm::DenseMap<const ValueDecl *, FieldDecl *> Captures;
2328 FieldDecl *ThisCapture = nullptr;
2329 RD->getCaptureFields(Captures, ThisCapture);
2330 if (ThisCapture && CGF.CapturedStmtInfo->isCXXThisExprCaptured()) {
2331 LValue ThisLVal =
2332 CGF.EmitLValueForFieldInitialization(Base: VDLVal, Field: ThisCapture);
2333 llvm::Value *CXXThis = CGF.LoadCXXThis();
2334 CGF.EmitStoreOfScalar(value: CXXThis, lvalue: ThisLVal);
2335 }
2336 for (const LambdaCapture &LC : RD->captures()) {
2337 if (LC.getCaptureKind() != LCK_ByRef)
2338 continue;
2339 const ValueDecl *VD = LC.getCapturedVar();
2340 // FIXME: For now VD is always a VarDecl because OpenMP does not support
2341 // capturing structured bindings in lambdas yet.
2342 if (!CS->capturesVariable(Var: cast<VarDecl>(Val: VD)))
2343 continue;
2344 auto It = Captures.find(Val: VD);
2345 assert(It != Captures.end() && "Found lambda capture without field.");
2346 LValue VarLVal = CGF.EmitLValueForFieldInitialization(Base: VDLVal, Field: It->second);
2347 Address VDAddr = CGF.GetAddrOfLocalVar(VD: cast<VarDecl>(Val: VD));
2348 if (VD->getType().getCanonicalType()->isReferenceType())
2349 VDAddr = CGF.EmitLoadOfReferenceLValue(RefAddr: VDAddr,
2350 RefTy: VD->getType().getCanonicalType())
2351 .getAddress();
2352 CGF.EmitStoreOfScalar(value: VDAddr.emitRawPointer(CGF), lvalue: VarLVal);
2353 }
2354 }
2355}
2356
2357bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
2358 LangAS &AS) {
2359 if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
2360 return false;
2361 const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
2362 switch(A->getAllocatorType()) {
2363 case OMPAllocateDeclAttr::OMPNullMemAlloc:
2364 case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
2365 // Not supported, fallback to the default mem space.
2366 case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
2367 case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
2368 case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
2369 case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
2370 case OMPAllocateDeclAttr::OMPThreadMemAlloc:
2371 AS = LangAS::Default;
2372 return true;
2373 case OMPAllocateDeclAttr::OMPConstMemAlloc:
2374 AS = LangAS::cuda_constant;
2375 return true;
2376 case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
2377 AS = LangAS::cuda_shared;
2378 return true;
2379 case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
2380 llvm_unreachable("Expected predefined allocator for the variables with the "
2381 "static storage.");
2382 }
2383 return false;
2384}
2385
2386/// Check to see if target architecture supports unified addressing which is
2387/// a restriction for OpenMP requires clause "unified_shared_memory".
2388void CGOpenMPRuntimeGPU::processRequiresDirective(const OMPRequiresDecl *D) {
2389 StringRef CPU = CGM.getTarget().getTargetOpts().CPU;
2390 if (CGM.getTarget().getTriple().isNVPTX() &&
2391 !llvm::NVPTX::supportsUnifiedAddressing(Kind: llvm::NVPTX::parseArch(CPU))) {
2392 for (const OMPClause *Clause : D->clauselists()) {
2393 if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
2394 CGM.getDiags().Report(Loc: Clause->getBeginLoc(),
2395 DiagID: diag::err_omp_unified_shared_memory_unsupported)
2396 << CPU;
2397 return;
2398 }
2399 }
2400 }
2401
2402 CGOpenMPRuntime::processRequiresDirective(D);
2403}
2404
2405llvm::Value *CGOpenMPRuntimeGPU::getGPUNumThreads(CodeGenFunction &CGF) {
2406 CGBuilderTy &Bld = CGF.Builder;
2407 llvm::Module *M = &CGF.CGM.getModule();
2408 const char *LocSize = "__kmpc_get_hardware_num_threads_in_block";
2409 llvm::Function *F = M->getFunction(Name: LocSize);
2410 if (!F) {
2411 F = llvm::Function::Create(Ty: llvm::FunctionType::get(Result: CGF.Int32Ty, Params: {}, isVarArg: false),
2412 Linkage: llvm::GlobalVariable::ExternalLinkage, N: LocSize,
2413 M: &CGF.CGM.getModule());
2414 }
2415 return Bld.CreateCall(Callee: F, Args: {}, Name: "nvptx_num_threads");
2416}
2417
2418llvm::Value *CGOpenMPRuntimeGPU::getGPUThreadID(CodeGenFunction &CGF) {
2419 ArrayRef<llvm::Value *> Args{};
2420 return CGF.EmitRuntimeCall(
2421 callee: OMPBuilder.getOrCreateRuntimeFunction(
2422 M&: CGM.getModule(), FnID: OMPRTL___kmpc_get_hardware_thread_id_in_block),
2423 args: Args);
2424}
2425