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::emitKernelInit(const OMPExecutableDirective &D,
751 CodeGenFunction &CGF,
752 EntryFunctionState &EST, bool IsSPMD) {
753 llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;
754 Attrs.ExecFlags =
755 IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD
756 : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;
757 computeMinAndMaxThreadsAndTeams(D, CGF, Attrs);
758
759 CGBuilderTy &Bld = CGF.Builder;
760 Bld.restoreIP(IP: OMPBuilder.createTargetInit(Loc: Bld, Attrs));
761 if (!IsSPMD)
762 emitGenericVarsProlog(CGF, Loc: EST.Loc);
763}
764
765void CGOpenMPRuntimeGPU::emitKernelDeinit(CodeGenFunction &CGF,
766 EntryFunctionState &EST,
767 bool IsSPMD) {
768 if (!IsSPMD)
769 emitGenericVarsEpilog(CGF);
770
771 // This is temporary until we remove the fixed sized buffer.
772 ASTContext &C = CGM.getContext();
773 RecordDecl *StaticRD = C.buildImplicitRecord(
774 Name: "_openmp_teams_reduction_type_$_", TK: RecordDecl::TagKind::Union);
775 StaticRD->startDefinition();
776 for (const RecordDecl *TeamReductionRec : TeamsReductions) {
777 CanQualType RecTy = C.getCanonicalTagType(TD: TeamReductionRec);
778 auto *Field = FieldDecl::Create(
779 C, DC: StaticRD, StartLoc: SourceLocation(), IdLoc: SourceLocation(), Id: nullptr, T: RecTy,
780 TInfo: C.getTrivialTypeSourceInfo(T: RecTy, Loc: SourceLocation()),
781 /*BW=*/nullptr, /*Mutable=*/false,
782 /*InitStyle=*/ICIS_NoInit);
783 Field->setAccess(AS_public);
784 StaticRD->addDecl(D: Field);
785 }
786 StaticRD->completeDefinition();
787 CanQualType StaticTy = C.getCanonicalTagType(TD: StaticRD);
788 llvm::Type *LLVMReductionsBufferTy =
789 CGM.getTypes().ConvertTypeForMem(T: StaticTy);
790 const auto &DL = CGM.getModule().getDataLayout();
791 uint64_t ReductionDataSize =
792 TeamsReductions.empty()
793 ? 0
794 : DL.getTypeAllocSize(Ty: LLVMReductionsBufferTy).getFixedValue();
795 CGBuilderTy &Bld = CGF.Builder;
796 OMPBuilder.createTargetDeinit(Loc: Bld, TeamsReductionDataSize: ReductionDataSize);
797 TeamsReductions.clear();
798}
799
800void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
801 StringRef ParentName,
802 llvm::Function *&OutlinedFn,
803 llvm::Constant *&OutlinedFnID,
804 bool IsOffloadEntry,
805 const RegionCodeGenTy &CodeGen) {
806 ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_SPMD);
807 EntryFunctionState EST;
808
809 bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
810
811 // Emit target region as a standalone region.
812 class NVPTXPrePostActionTy : public PrePostActionTy {
813 CGOpenMPRuntimeGPU &RT;
814 CGOpenMPRuntimeGPU::EntryFunctionState &EST;
815 bool IsBareKernel;
816 DataSharingMode Mode;
817 const OMPExecutableDirective &D;
818
819 public:
820 NVPTXPrePostActionTy(CGOpenMPRuntimeGPU &RT,
821 CGOpenMPRuntimeGPU::EntryFunctionState &EST,
822 bool IsBareKernel, const OMPExecutableDirective &D)
823 : RT(RT), EST(EST), IsBareKernel(IsBareKernel),
824 Mode(RT.CurrentDataSharingMode), D(D) {}
825 void Enter(CodeGenFunction &CGF) override {
826 if (IsBareKernel) {
827 RT.CurrentDataSharingMode = DataSharingMode::DS_CUDA;
828 return;
829 }
830 RT.emitKernelInit(D, CGF, EST, /* IsSPMD */ true);
831 // Skip target region initialization.
832 RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
833 }
834 void Exit(CodeGenFunction &CGF) override {
835 if (IsBareKernel) {
836 RT.CurrentDataSharingMode = Mode;
837 return;
838 }
839 RT.clearLocThreadIdInsertPt(CGF);
840 RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ true);
841 }
842 } Action(*this, EST, IsBareKernel, D);
843 CodeGen.setAction(Action);
844 IsInTTDRegion = true;
845 emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
846 IsOffloadEntry, CodeGen);
847 IsInTTDRegion = false;
848}
849
850void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
851 const OMPExecutableDirective &D, StringRef ParentName,
852 llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
853 bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
854 if (!IsOffloadEntry) // Nothing to do.
855 return;
856
857 assert(!ParentName.empty() && "Invalid target region parent name!");
858
859 bool Mode = supportsSPMDExecutionMode(Ctx&: CGM.getContext(), D);
860 bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
861 if (Mode || IsBareKernel)
862 emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
863 CodeGen);
864 else
865 emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
866 CodeGen);
867}
868
869CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
870 : CGOpenMPRuntime(CGM) {
871 llvm::OpenMPIRBuilderConfig Config(
872 CGM.getLangOpts().OpenMPIsTargetDevice, isGPU(),
873 CGM.getLangOpts().OpenMPOffloadMandatory,
874 /*HasRequiresReverseOffload*/ false, /*HasRequiresUnifiedAddress*/ false,
875 hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
876 Config.setDefaultTargetAS(
877 CGM.getContext().getTargetInfo().getTargetAddressSpace(AS: LangAS::Default));
878 Config.setRuntimeCC(CGM.getRuntimeCC());
879
880 OMPBuilder.setConfig(Config);
881
882 if (!CGM.getLangOpts().OpenMPIsTargetDevice)
883 llvm_unreachable("OpenMP can only handle device code.");
884
885 if (CGM.getLangOpts().OpenMPCUDAMode)
886 CurrentDataSharingMode = CGOpenMPRuntimeGPU::DS_CUDA;
887
888 llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder();
889 if (CGM.getLangOpts().NoGPULib || CGM.getLangOpts().OMPHostIRFile.empty())
890 return;
891
892 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPTargetDebug,
893 Name: "__omp_rtl_debug_kind");
894 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPTeamSubscription,
895 Name: "__omp_rtl_assume_teams_oversubscription");
896 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPThreadSubscription,
897 Name: "__omp_rtl_assume_threads_oversubscription");
898 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPNoThreadState,
899 Name: "__omp_rtl_assume_no_thread_state");
900 OMPBuilder.createGlobalFlag(Value: CGM.getLangOpts().OpenMPNoNestedParallelism,
901 Name: "__omp_rtl_assume_no_nested_parallelism");
902}
903
904void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,
905 ProcBindKind ProcBind,
906 SourceLocation Loc) {
907 // Nothing to do.
908}
909
910llvm::Value *CGOpenMPRuntimeGPU::emitMessageClause(CodeGenFunction &CGF,
911 const Expr *Message,
912 SourceLocation Loc) {
913 CGM.getDiags().Report(Loc, DiagID: diag::warn_omp_gpu_unsupported_clause)
914 << getOpenMPClauseName(C: OMPC_message);
915 return nullptr;
916}
917
918llvm::Value *
919CGOpenMPRuntimeGPU::emitSeverityClause(OpenMPSeverityClauseKind Severity,
920 SourceLocation Loc) {
921 CGM.getDiags().Report(Loc, DiagID: diag::warn_omp_gpu_unsupported_clause)
922 << getOpenMPClauseName(C: OMPC_severity);
923 return nullptr;
924}
925
926void CGOpenMPRuntimeGPU::emitNumThreadsClause(
927 CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
928 OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,
929 SourceLocation SeverityLoc, const Expr *Message,
930 SourceLocation MessageLoc) {
931 if (Modifier == OMPC_NUMTHREADS_strict) {
932 CGM.getDiags().Report(Loc,
933 DiagID: diag::warn_omp_gpu_unsupported_modifier_for_clause)
934 << "strict" << getOpenMPClauseName(C: OMPC_num_threads);
935 return;
936 }
937
938 // Nothing to do.
939}
940
941void CGOpenMPRuntimeGPU::emitNumTeamsClause(CodeGenFunction &CGF,
942 const Expr *NumTeams,
943 const Expr *ThreadLimit,
944 SourceLocation Loc) {}
945
946llvm::Function *CGOpenMPRuntimeGPU::emitParallelOutlinedFunction(
947 CodeGenFunction &CGF, const OMPExecutableDirective &D,
948 const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
949 const RegionCodeGenTy &CodeGen) {
950 // Emit target region as a standalone region.
951 bool PrevIsInTTDRegion = IsInTTDRegion;
952 IsInTTDRegion = false;
953 auto *OutlinedFun =
954 cast<llvm::Function>(Val: CGOpenMPRuntime::emitParallelOutlinedFunction(
955 CGF, D, ThreadIDVar, InnermostKind, CodeGen));
956 IsInTTDRegion = PrevIsInTTDRegion;
957 if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD) {
958 llvm::Function *WrapperFun =
959 createParallelDataSharingWrapper(OutlinedParallelFn: OutlinedFun, D);
960 WrapperFunctionsMap[OutlinedFun] = WrapperFun;
961 }
962
963 return OutlinedFun;
964}
965
966/// Get list of lastprivate variables from the teams distribute ... or
967/// teams {distribute ...} directives.
968static void
969getDistributeLastprivateVars(ASTContext &Ctx, const OMPExecutableDirective &D,
970 llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
971 assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
972 "expected teams directive.");
973 const OMPExecutableDirective *Dir = &D;
974 if (!isOpenMPDistributeDirective(DKind: D.getDirectiveKind())) {
975 if (const Stmt *S = CGOpenMPRuntime::getSingleCompoundChild(
976 Ctx,
977 Body: D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(
978 /*IgnoreCaptured=*/true))) {
979 Dir = dyn_cast_or_null<OMPExecutableDirective>(Val: S);
980 if (Dir && !isOpenMPDistributeDirective(DKind: Dir->getDirectiveKind()))
981 Dir = nullptr;
982 }
983 }
984 if (!Dir)
985 return;
986 for (const auto *C : Dir->getClausesOfKind<OMPLastprivateClause>()) {
987 for (const Expr *E : C->getVarRefs())
988 Vars.push_back(Elt: getPrivateItem(RefExpr: E));
989 }
990}
991
992/// Get list of reduction variables from the teams ... directives.
993static void
994getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D,
995 llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
996 assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
997 "expected teams directive.");
998 for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
999 for (const Expr *E : C->privates())
1000 Vars.push_back(Elt: getPrivateItem(RefExpr: E));
1001 }
1002}
1003
1004llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(
1005 CodeGenFunction &CGF, const OMPExecutableDirective &D,
1006 const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
1007 const RegionCodeGenTy &CodeGen) {
1008 SourceLocation Loc = D.getBeginLoc();
1009
1010 const RecordDecl *GlobalizedRD = nullptr;
1011 llvm::SmallVector<const ValueDecl *, 4> LastPrivatesReductions;
1012 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
1013 unsigned WarpSize = CGM.getTarget().getGridValue().GV_Warp_Size;
1014 // Globalize team reductions variable unconditionally in all modes.
1015 if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1016 getTeamsReductionVars(Ctx&: CGM.getContext(), D, Vars&: LastPrivatesReductions);
1017 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
1018 getDistributeLastprivateVars(Ctx&: CGM.getContext(), D, Vars&: LastPrivatesReductions);
1019 if (!LastPrivatesReductions.empty()) {
1020 GlobalizedRD = ::buildRecordForGlobalizedVars(
1021 C&: CGM.getContext(), EscapedDecls: {}, EscapedDeclsForTeams: LastPrivatesReductions, MappedDeclsFields,
1022 BufSize: WarpSize);
1023 }
1024 } else if (!LastPrivatesReductions.empty()) {
1025 assert(!TeamAndReductions.first &&
1026 "Previous team declaration is not expected.");
1027 TeamAndReductions.first = D.getCapturedStmt(RegionKind: OMPD_teams)->getCapturedDecl();
1028 std::swap(LHS&: TeamAndReductions.second, RHS&: LastPrivatesReductions);
1029 }
1030
1031 // Emit target region as a standalone region.
1032 class NVPTXPrePostActionTy : public PrePostActionTy {
1033 SourceLocation &Loc;
1034 const RecordDecl *GlobalizedRD;
1035 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1036 &MappedDeclsFields;
1037
1038 public:
1039 NVPTXPrePostActionTy(
1040 SourceLocation &Loc, const RecordDecl *GlobalizedRD,
1041 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1042 &MappedDeclsFields)
1043 : Loc(Loc), GlobalizedRD(GlobalizedRD),
1044 MappedDeclsFields(MappedDeclsFields) {}
1045 void Enter(CodeGenFunction &CGF) override {
1046 auto &Rt =
1047 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1048 if (GlobalizedRD) {
1049 auto I = Rt.FunctionGlobalizedDecls.try_emplace(Key: CGF.CurFn).first;
1050 I->getSecond().MappedParams =
1051 std::make_unique<CodeGenFunction::OMPMapVars>();
1052 DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
1053 for (const auto &Pair : MappedDeclsFields) {
1054 assert(Pair.getFirst()->isCanonicalDecl() &&
1055 "Expected canonical declaration");
1056 Data.try_emplace(Key: Pair.getFirst());
1057 }
1058 }
1059 Rt.emitGenericVarsProlog(CGF, Loc);
1060 }
1061 void Exit(CodeGenFunction &CGF) override {
1062 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
1063 .emitGenericVarsEpilog(CGF);
1064 }
1065 } Action(Loc, GlobalizedRD, MappedDeclsFields);
1066 CodeGen.setAction(Action);
1067 llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(
1068 CGF, D, ThreadIDVar, InnermostKind, CodeGen);
1069
1070 return OutlinedFun;
1071}
1072
1073void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,
1074 SourceLocation Loc) {
1075 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
1076 return;
1077
1078 CGBuilderTy &Bld = CGF.Builder;
1079
1080 const auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
1081 if (I == FunctionGlobalizedDecls.end())
1082 return;
1083
1084 for (auto &Rec : I->getSecond().LocalVarData) {
1085 const auto *VD = cast<VarDecl>(Val: Rec.first);
1086 bool EscapedParam = I->getSecond().EscapedParameters.count(Ptr: Rec.first);
1087 QualType VarTy = VD->getType();
1088
1089 // Get the local allocation of a firstprivate variable before sharing
1090 llvm::Value *ParValue;
1091 if (EscapedParam) {
1092 LValue ParLVal =
1093 CGF.MakeAddrLValue(Addr: CGF.GetAddrOfLocalVar(VD), T: VD->getType());
1094 ParValue = CGF.EmitLoadOfScalar(lvalue: ParLVal, Loc);
1095 }
1096
1097 // Allocate space for the variable to be globalized
1098 llvm::Value *AllocArgs[] = {CGF.getTypeSize(Ty: VD->getType())};
1099 llvm::CallBase *VoidPtr =
1100 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1101 M&: CGM.getModule(), FnID: OMPRTL___kmpc_alloc_shared),
1102 args: AllocArgs, name: VD->getName());
1103 // FIXME: We should use the variables actual alignment as an argument.
1104 VoidPtr->addRetAttr(Attr: llvm::Attribute::get(
1105 Context&: CGM.getLLVMContext(), Kind: llvm::Attribute::Alignment,
1106 Val: CGM.getContext().getTargetInfo().getNewAlign() / 8));
1107
1108 // Cast the void pointer and get the address of the globalized variable.
1109 llvm::Value *CastedVoidPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1110 V: VoidPtr, DestTy: Bld.getPtrTy(AddrSpace: 0), Name: VD->getName() + "_on_stack");
1111 LValue VarAddr =
1112 CGF.MakeNaturalAlignPointeeRawAddrLValue(V: CastedVoidPtr, T: VarTy);
1113 Rec.second.PrivateAddr = VarAddr.getAddress();
1114 Rec.second.GlobalizedVal = VoidPtr;
1115
1116 // Assign the local allocation to the newly globalized location.
1117 if (EscapedParam) {
1118 CGF.EmitStoreOfScalar(value: ParValue, lvalue: VarAddr);
1119 I->getSecond().MappedParams->setVarAddr(CGF, LocalVD: VD, TempAddr: VarAddr.getAddress());
1120 }
1121 if (auto *DI = CGF.getDebugInfo())
1122 VoidPtr->setDebugLoc(DI->SourceLocToDebugLoc(Loc: VD->getLocation()));
1123 }
1124
1125 for (const auto *ValueD : I->getSecond().EscapedVariableLengthDecls) {
1126 const auto *VD = cast<VarDecl>(Val: ValueD);
1127 std::pair<llvm::Value *, llvm::Value *> AddrSizePair =
1128 getKmpcAllocShared(CGF, VD);
1129 I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(Args&: AddrSizePair);
1130 LValue Base = CGF.MakeAddrLValue(V: AddrSizePair.first, T: VD->getType(),
1131 Alignment: CGM.getContext().getDeclAlign(D: VD),
1132 Source: AlignmentSource::Decl);
1133 I->getSecond().MappedParams->setVarAddr(CGF, LocalVD: VD, TempAddr: Base.getAddress());
1134 }
1135 I->getSecond().MappedParams->apply(CGF);
1136}
1137
1138bool CGOpenMPRuntimeGPU::isDelayedVariableLengthDecl(CodeGenFunction &CGF,
1139 const VarDecl *VD) const {
1140 const auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
1141 if (I == FunctionGlobalizedDecls.end())
1142 return false;
1143
1144 // Check variable declaration is delayed:
1145 return llvm::is_contained(Range: I->getSecond().DelayedVariableLengthDecls, Element: VD);
1146}
1147
1148std::pair<llvm::Value *, llvm::Value *>
1149CGOpenMPRuntimeGPU::getKmpcAllocShared(CodeGenFunction &CGF,
1150 const VarDecl *VD) {
1151 CGBuilderTy &Bld = CGF.Builder;
1152
1153 // Compute size and alignment.
1154 llvm::Value *Size = CGF.getTypeSize(Ty: VD->getType());
1155 CharUnits Align = CGM.getContext().getDeclAlign(D: VD);
1156 Size = Bld.CreateNUWAdd(
1157 LHS: Size, RHS: llvm::ConstantInt::get(Ty: CGF.SizeTy, V: Align.getQuantity() - 1));
1158 llvm::Value *AlignVal =
1159 llvm::ConstantInt::get(Ty: CGF.SizeTy, V: Align.getQuantity());
1160 Size = Bld.CreateUDiv(LHS: Size, RHS: AlignVal);
1161 Size = Bld.CreateNUWMul(LHS: Size, RHS: AlignVal);
1162
1163 // Allocate space for this VLA object to be globalized.
1164 llvm::Value *AllocArgs[] = {Size};
1165 llvm::CallBase *VoidPtr =
1166 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1167 M&: CGM.getModule(), FnID: OMPRTL___kmpc_alloc_shared),
1168 args: AllocArgs, name: VD->getName());
1169 VoidPtr->addRetAttr(Attr: llvm::Attribute::get(
1170 Context&: CGM.getLLVMContext(), Kind: llvm::Attribute::Alignment, Val: Align.getQuantity()));
1171
1172 return std::make_pair(x&: VoidPtr, y&: Size);
1173}
1174
1175void CGOpenMPRuntimeGPU::getKmpcFreeShared(
1176 CodeGenFunction &CGF,
1177 const std::pair<llvm::Value *, llvm::Value *> &AddrSizePair) {
1178 // Deallocate the memory for each globalized VLA object
1179 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1180 M&: CGM.getModule(), FnID: OMPRTL___kmpc_free_shared),
1181 args: {AddrSizePair.first, AddrSizePair.second});
1182}
1183
1184void CGOpenMPRuntimeGPU::emitGenericVarsEpilog(CodeGenFunction &CGF) {
1185 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
1186 return;
1187
1188 const auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
1189 if (I != FunctionGlobalizedDecls.end()) {
1190 // Deallocate the memory for each globalized VLA object that was
1191 // globalized in the prolog (i.e. emitGenericVarsProlog).
1192 for (const auto &AddrSizePair :
1193 llvm::reverse(C&: I->getSecond().EscapedVariableLengthDeclsAddrs)) {
1194 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1195 M&: CGM.getModule(), FnID: OMPRTL___kmpc_free_shared),
1196 args: {AddrSizePair.first, AddrSizePair.second});
1197 }
1198 // Deallocate the memory for each globalized value
1199 for (auto &Rec : llvm::reverse(C&: I->getSecond().LocalVarData)) {
1200 const auto *VD = cast<VarDecl>(Val: Rec.first);
1201 I->getSecond().MappedParams->restore(CGF);
1202
1203 llvm::Value *FreeArgs[] = {Rec.second.GlobalizedVal,
1204 CGF.getTypeSize(Ty: VD->getType())};
1205 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1206 M&: CGM.getModule(), FnID: OMPRTL___kmpc_free_shared),
1207 args: FreeArgs);
1208 }
1209 }
1210}
1211
1212void CGOpenMPRuntimeGPU::emitTeamsCall(CodeGenFunction &CGF,
1213 const OMPExecutableDirective &D,
1214 SourceLocation Loc,
1215 llvm::Function *OutlinedFn,
1216 ArrayRef<llvm::Value *> CapturedVars) {
1217 if (!CGF.HaveInsertPoint())
1218 return;
1219
1220 bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
1221
1222 RawAddress ZeroAddr = CGF.CreateDefaultAlignTempAlloca(Ty: CGF.Int32Ty,
1223 /*Name=*/".zero.addr");
1224 CGF.Builder.CreateStore(Val: CGF.Builder.getInt32(/*C*/ 0), Addr: ZeroAddr);
1225 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
1226 // We don't emit any thread id function call in bare kernel, but because the
1227 // outlined function has a pointer argument, we emit a nullptr here.
1228 if (IsBareKernel)
1229 OutlinedFnArgs.push_back(Elt: llvm::ConstantPointerNull::get(T: CGM.VoidPtrTy));
1230 else
1231 OutlinedFnArgs.push_back(Elt: emitThreadIDAddress(CGF, Loc).emitRawPointer(CGF));
1232 OutlinedFnArgs.push_back(Elt: ZeroAddr.getPointer());
1233 OutlinedFnArgs.append(in_start: CapturedVars.begin(), in_end: CapturedVars.end());
1234 emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, Args: OutlinedFnArgs);
1235}
1236
1237void CGOpenMPRuntimeGPU::emitParallelCall(
1238 CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *OutlinedFn,
1239 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond,
1240 llvm::Value *NumThreads, OpenMPNumThreadsClauseModifier NumThreadsModifier,
1241 OpenMPSeverityClauseKind Severity, const Expr *Message) {
1242 if (!CGF.HaveInsertPoint())
1243 return;
1244
1245 auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
1246 NumThreads](CodeGenFunction &CGF,
1247 PrePostActionTy &Action) {
1248 CGBuilderTy &Bld = CGF.Builder;
1249 llvm::Value *NumThreadsVal = NumThreads;
1250 llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
1251 llvm::PointerType *FnPtrTy = llvm::PointerType::get(
1252 C&: CGF.getLLVMContext(), AddressSpace: CGM.getDataLayout().getProgramAddressSpace());
1253
1254 llvm::Value *ID = llvm::ConstantPointerNull::get(T: FnPtrTy);
1255 if (WFn)
1256 ID = Bld.CreateBitOrPointerCast(V: WFn, DestTy: FnPtrTy);
1257
1258 llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(V: OutlinedFn, DestTy: FnPtrTy);
1259
1260 // Create a private scope that will globalize the arguments
1261 // passed from the outside of the target region.
1262 // TODO: Is that needed?
1263 CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF);
1264
1265 Address CapturedVarsAddrs = CGF.CreateDefaultAlignTempAlloca(
1266 Ty: llvm::ArrayType::get(ElementType: CGM.VoidPtrTy, NumElements: CapturedVars.size()),
1267 Name: "captured_vars_addrs");
1268 // There's something to share.
1269 if (!CapturedVars.empty()) {
1270 // Prepare for parallel region. Indicate the outlined function.
1271 ASTContext &Ctx = CGF.getContext();
1272 unsigned Idx = 0;
1273 for (llvm::Value *V : CapturedVars) {
1274 Address Dst = Bld.CreateConstArrayGEP(Addr: CapturedVarsAddrs, Index: Idx);
1275 llvm::Value *PtrV;
1276 if (V->getType()->isIntegerTy())
1277 PtrV = Bld.CreateIntToPtr(V, DestTy: CGF.VoidPtrTy);
1278 else
1279 PtrV = Bld.CreatePointerBitCastOrAddrSpaceCast(V, DestTy: CGF.VoidPtrTy);
1280 CGF.EmitStoreOfScalar(Value: PtrV, Addr: Dst, /*Volatile=*/false,
1281 Ty: Ctx.getPointerType(T: Ctx.VoidPtrTy));
1282 ++Idx;
1283 }
1284 }
1285
1286 llvm::Value *IfCondVal = nullptr;
1287 if (IfCond)
1288 IfCondVal = Bld.CreateIntCast(V: CGF.EvaluateExprAsBool(E: IfCond), DestTy: CGF.Int32Ty,
1289 /* isSigned */ false);
1290 else
1291 IfCondVal = llvm::ConstantInt::get(Ty: CGF.Int32Ty, V: 1);
1292
1293 if (!NumThreadsVal)
1294 NumThreadsVal = llvm::ConstantInt::getAllOnesValue(Ty: CGF.Int32Ty);
1295 else
1296 NumThreadsVal = Bld.CreateZExtOrTrunc(V: NumThreadsVal, DestTy: CGF.Int32Ty);
1297
1298 // No strict prescriptiveness for the number of threads.
1299 llvm::Value *StrictNumThreadsVal = llvm::ConstantInt::get(Ty: CGF.Int32Ty, V: 0);
1300
1301 assert(IfCondVal && "Expected a value");
1302 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
1303 llvm::Value *Args[] = {
1304 RTLoc,
1305 getThreadID(CGF, Loc),
1306 IfCondVal,
1307 NumThreadsVal,
1308 llvm::ConstantInt::getAllOnesValue(Ty: CGF.Int32Ty),
1309 FnPtr,
1310 ID,
1311 Bld.CreateBitOrPointerCast(V: CapturedVarsAddrs.emitRawPointer(CGF),
1312 DestTy: CGF.VoidPtrPtrTy),
1313 llvm::ConstantInt::get(Ty: CGM.SizeTy, V: CapturedVars.size()),
1314 StrictNumThreadsVal};
1315
1316 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1317 M&: CGM.getModule(), FnID: OMPRTL___kmpc_parallel_60),
1318 args: Args);
1319 };
1320
1321 RegionCodeGenTy RCG(ParallelGen);
1322 RCG(CGF);
1323}
1324
1325void CGOpenMPRuntimeGPU::syncCTAThreads(CodeGenFunction &CGF) {
1326 // Always emit simple barriers!
1327 if (!CGF.HaveInsertPoint())
1328 return;
1329 // Build call __kmpc_barrier_simple_spmd(nullptr, 0);
1330 // This function does not use parameters, so we can emit just default values.
1331 llvm::Value *Args[] = {
1332 llvm::ConstantPointerNull::get(
1333 T: cast<llvm::PointerType>(Val: getIdentTyPointerTy())),
1334 llvm::ConstantInt::get(Ty: CGF.Int32Ty, /*V=*/0, /*isSigned=*/IsSigned: true)};
1335 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1336 M&: CGM.getModule(), FnID: OMPRTL___kmpc_barrier_simple_spmd),
1337 args: Args);
1338}
1339
1340void CGOpenMPRuntimeGPU::emitBarrierCall(CodeGenFunction &CGF,
1341 SourceLocation Loc,
1342 OpenMPDirectiveKind Kind, bool,
1343 bool) {
1344 // Always emit simple barriers!
1345 if (!CGF.HaveInsertPoint())
1346 return;
1347 // Build call __kmpc_cancel_barrier(loc, thread_id);
1348 unsigned Flags = getDefaultFlagsForBarriers(Kind);
1349 llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc, Flags),
1350 getThreadID(CGF, Loc)};
1351
1352 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1353 M&: CGM.getModule(), FnID: OMPRTL___kmpc_barrier),
1354 args: Args);
1355}
1356
1357void CGOpenMPRuntimeGPU::emitCriticalRegion(
1358 CodeGenFunction &CGF, StringRef CriticalName,
1359 const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc,
1360 const Expr *Hint) {
1361 llvm::BasicBlock *LoopBB = CGF.createBasicBlock(name: "omp.critical.loop");
1362 llvm::BasicBlock *TestBB = CGF.createBasicBlock(name: "omp.critical.test");
1363 llvm::BasicBlock *SyncBB = CGF.createBasicBlock(name: "omp.critical.sync");
1364 llvm::BasicBlock *BodyBB = CGF.createBasicBlock(name: "omp.critical.body");
1365 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(name: "omp.critical.exit");
1366
1367 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1368
1369 // Get the mask of active threads in the warp.
1370 llvm::Value *Mask = CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1371 M&: CGM.getModule(), FnID: OMPRTL___kmpc_warp_active_thread_mask));
1372 // Fetch team-local id of the thread.
1373 llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
1374
1375 // Get the width of the team.
1376 llvm::Value *TeamWidth = RT.getGPUNumThreads(CGF);
1377
1378 // Initialize the counter variable for the loop.
1379 QualType Int32Ty =
1380 CGF.getContext().getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/0);
1381 Address Counter = CGF.CreateMemTempWithoutCast(T: Int32Ty, Name: "critical_counter");
1382 LValue CounterLVal = CGF.MakeAddrLValue(Addr: Counter, T: Int32Ty);
1383 CGF.EmitStoreOfScalar(value: llvm::Constant::getNullValue(Ty: CGM.Int32Ty), lvalue: CounterLVal,
1384 /*isInit=*/true);
1385
1386 // Block checks if loop counter exceeds upper bound.
1387 CGF.EmitBlock(BB: LoopBB);
1388 llvm::Value *CounterVal = CGF.EmitLoadOfScalar(lvalue: CounterLVal, Loc);
1389 llvm::Value *CmpLoopBound = CGF.Builder.CreateICmpSLT(LHS: CounterVal, RHS: TeamWidth);
1390 CGF.Builder.CreateCondBr(Cond: CmpLoopBound, True: TestBB, False: ExitBB);
1391
1392 // Block tests which single thread should execute region, and which threads
1393 // should go straight to synchronisation point.
1394 CGF.EmitBlock(BB: TestBB);
1395 CounterVal = CGF.EmitLoadOfScalar(lvalue: CounterLVal, Loc);
1396 llvm::Value *CmpThreadToCounter =
1397 CGF.Builder.CreateICmpEQ(LHS: ThreadID, RHS: CounterVal);
1398 CGF.Builder.CreateCondBr(Cond: CmpThreadToCounter, True: BodyBB, False: SyncBB);
1399
1400 // Block emits the body of the critical region.
1401 CGF.EmitBlock(BB: BodyBB);
1402
1403 // Output the critical statement.
1404 CGOpenMPRuntime::emitCriticalRegion(CGF, CriticalName, CriticalOpGen, Loc,
1405 Hint);
1406
1407 // After the body surrounded by the critical region, the single executing
1408 // thread will jump to the synchronisation point.
1409 // Block waits for all threads in current team to finish then increments the
1410 // counter variable and returns to the loop.
1411 CGF.EmitBlock(BB: SyncBB);
1412 // Reconverge active threads in the warp.
1413 (void)CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
1414 M&: CGM.getModule(), FnID: OMPRTL___kmpc_syncwarp),
1415 args: Mask);
1416
1417 llvm::Value *IncCounterVal =
1418 CGF.Builder.CreateNSWAdd(LHS: CounterVal, RHS: CGF.Builder.getInt32(C: 1));
1419 CGF.EmitStoreOfScalar(value: IncCounterVal, lvalue: CounterLVal);
1420 CGF.EmitBranch(Block: LoopBB);
1421
1422 // Block that is reached when all threads in the team complete the region.
1423 CGF.EmitBlock(BB: ExitBB, /*IsFinished=*/true);
1424}
1425
1426/// Cast value to the specified type.
1427static llvm::Value *castValueToType(CodeGenFunction &CGF, llvm::Value *Val,
1428 QualType ValTy, QualType CastTy,
1429 SourceLocation Loc) {
1430 assert(!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&
1431 "Cast type must sized.");
1432 assert(!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&
1433 "Val type must sized.");
1434 llvm::Type *LLVMCastTy = CGF.ConvertTypeForMem(T: CastTy);
1435 if (ValTy == CastTy)
1436 return Val;
1437 if (CGF.getContext().getTypeSizeInChars(T: ValTy) ==
1438 CGF.getContext().getTypeSizeInChars(T: CastTy))
1439 return CGF.Builder.CreateBitCast(V: Val, DestTy: LLVMCastTy);
1440 if (CastTy->isIntegerType() && ValTy->isIntegerType())
1441 return CGF.Builder.CreateIntCast(V: Val, DestTy: LLVMCastTy,
1442 isSigned: CastTy->hasSignedIntegerRepresentation());
1443 Address CastItem = CGF.CreateMemTempWithoutCast(T: CastTy);
1444 Address ValCastItem = CastItem.withElementType(ElemTy: Val->getType());
1445 CGF.EmitStoreOfScalar(Value: Val, Addr: ValCastItem, /*Volatile=*/false, Ty: ValTy,
1446 BaseInfo: LValueBaseInfo(AlignmentSource::Type),
1447 TBAAInfo: TBAAAccessInfo());
1448 return CGF.EmitLoadOfScalar(Addr: CastItem, /*Volatile=*/false, Ty: CastTy, Loc,
1449 BaseInfo: LValueBaseInfo(AlignmentSource::Type),
1450 TBAAInfo: TBAAAccessInfo());
1451}
1452
1453/// Extracts the built-in reduction operator from a combiner of the form `x = x
1454/// <op> rhs` (or the min/max conditional), or nullopt if the shape is not
1455/// recognized (e.g. user-defined reductions).
1456static std::optional<BinaryOperatorKind>
1457getReductionBinOpKind(const Expr *ReductionOp) {
1458 const auto *Assign = dyn_cast<BinaryOperator>(Val: ReductionOp);
1459 if (!Assign || Assign->getOpcode() != BO_Assign)
1460 return std::nullopt;
1461 const Expr *RHS = Assign->getRHS();
1462 // min/max are lowered as `x <cmp> rhs ? x : rhs`; the comparison identifies
1463 // it.
1464 if (const auto *ACO =
1465 dyn_cast<AbstractConditionalOperator>(Val: RHS->IgnoreParenImpCasts()))
1466 RHS = ACO->getCond();
1467 if (const auto *BO = dyn_cast<BinaryOperator>(Val: RHS->IgnoreParenImpCasts()))
1468 return BO->getOpcode();
1469 return std::nullopt;
1470}
1471
1472/// Maps a built-in reduction operator to an atomicrmw opcode for the atomic
1473/// cross-team reduction fast path, or nullopt if there is no direct atomicrmw
1474/// (e.g. user-defined, complex, fp min/max) so the buffer path is used instead.
1475static std::optional<llvm::AtomicRMWInst::BinOp>
1476getReductionAtomicRMWOp(BinaryOperatorKind BOK, QualType Ty) {
1477 bool IsInt = Ty->isIntegerType();
1478 bool IsSigned = Ty->hasSignedIntegerRepresentation();
1479 switch (BOK) {
1480 case BO_Add:
1481 case BO_Sub: // A `-` reduction sums the partials, so it accumulates with add.
1482 if (IsInt)
1483 return llvm::AtomicRMWInst::Add;
1484 if (Ty->isFloatingType())
1485 return llvm::AtomicRMWInst::FAdd;
1486 return std::nullopt;
1487 case BO_And:
1488 return IsInt ? std::optional(llvm::AtomicRMWInst::And) : std::nullopt;
1489 case BO_Or:
1490 return IsInt ? std::optional(llvm::AtomicRMWInst::Or) : std::nullopt;
1491 case BO_Xor:
1492 return IsInt ? std::optional(llvm::AtomicRMWInst::Xor) : std::nullopt;
1493 case BO_LT: // min
1494 if (IsInt)
1495 return IsSigned ? llvm::AtomicRMWInst::Min : llvm::AtomicRMWInst::UMin;
1496 return std::nullopt;
1497 case BO_GT: // max
1498 if (IsInt)
1499 return IsSigned ? llvm::AtomicRMWInst::Max : llvm::AtomicRMWInst::UMax;
1500 return std::nullopt;
1501 default:
1502 return std::nullopt;
1503 }
1504}
1505
1506///
1507/// Design of OpenMP reductions on the GPU
1508///
1509/// Consider a typical OpenMP program with one or more reduction
1510/// clauses:
1511///
1512/// float foo;
1513/// double bar;
1514/// #pragma omp target teams distribute parallel for \
1515/// reduction(+:foo) reduction(*:bar)
1516/// for (int i = 0; i < N; i++) {
1517/// foo += A[i]; bar *= B[i];
1518/// }
1519///
1520/// where 'foo' and 'bar' are reduced across all OpenMP threads in
1521/// all teams. In our OpenMP implementation on the NVPTX device an
1522/// OpenMP team is mapped to a CUDA threadblock and OpenMP threads
1523/// within a team are mapped to CUDA threads within a threadblock.
1524/// Our goal is to efficiently aggregate values across all OpenMP
1525/// threads such that:
1526///
1527/// - the compiler and runtime are logically concise, and
1528/// - the reduction is performed efficiently in a hierarchical
1529/// manner as follows: within OpenMP threads in the same warp,
1530/// across warps in a threadblock, and finally across teams on
1531/// the NVPTX device.
1532///
1533/// Introduction to Decoupling
1534///
1535/// We would like to decouple the compiler and the runtime so that the
1536/// latter is ignorant of the reduction variables (number, data types)
1537/// and the reduction operators. This allows a simpler interface
1538/// and implementation while still attaining good performance.
1539///
1540/// Pseudocode for the aforementioned OpenMP program generated by the
1541/// compiler is as follows:
1542///
1543/// 1. Create private copies of reduction variables on each OpenMP
1544/// thread: 'foo_private', 'bar_private'
1545/// 2. Each OpenMP thread reduces the chunk of 'A' and 'B' assigned
1546/// to it and writes the result in 'foo_private' and 'bar_private'
1547/// respectively.
1548/// 3. Call the OpenMP runtime on the GPU to reduce within a team
1549/// and store the result on the team master:
1550///
1551/// __kmpc_nvptx_parallel_reduce_nowait_v2(...,
1552/// reduceData, shuffleReduceFn, interWarpCpyFn)
1553///
1554/// where:
1555/// struct ReduceData {
1556/// double *foo;
1557/// double *bar;
1558/// } reduceData
1559/// reduceData.foo = &foo_private
1560/// reduceData.bar = &bar_private
1561///
1562/// 'shuffleReduceFn' and 'interWarpCpyFn' are pointers to two
1563/// auxiliary functions generated by the compiler that operate on
1564/// variables of type 'ReduceData'. They aid the runtime perform
1565/// algorithmic steps in a data agnostic manner.
1566///
1567/// 'shuffleReduceFn' is a pointer to a function that reduces data
1568/// of type 'ReduceData' across two OpenMP threads (lanes) in the
1569/// same warp. It takes the following arguments as input:
1570///
1571/// a. variable of type 'ReduceData' on the calling lane,
1572/// b. its lane_id,
1573/// c. an offset relative to the current lane_id to generate a
1574/// remote_lane_id. The remote lane contains the second
1575/// variable of type 'ReduceData' that is to be reduced.
1576/// d. an algorithm version parameter determining which reduction
1577/// algorithm to use.
1578///
1579/// 'shuffleReduceFn' retrieves data from the remote lane using
1580/// efficient GPU shuffle intrinsics and reduces, using the
1581/// algorithm specified by the 4th parameter, the two operands
1582/// element-wise. The result is written to the first operand.
1583///
1584/// Different reduction algorithms are implemented in different
1585/// runtime functions, all calling 'shuffleReduceFn' to perform
1586/// the essential reduction step. Therefore, based on the 4th
1587/// parameter, this function behaves slightly differently to
1588/// cooperate with the runtime to ensure correctness under
1589/// different circumstances.
1590///
1591/// 'InterWarpCpyFn' is a pointer to a function that transfers
1592/// reduced variables across warps. It tunnels, through CUDA
1593/// shared memory, the thread-private data of type 'ReduceData'
1594/// from lane 0 of each warp to a lane in the first warp.
1595/// 4. Call the OpenMP runtime on the GPU to reduce across teams.
1596/// The last team writes the global reduced value to memory.
1597///
1598/// ret = __kmpc_nvptx_teams_reduce_nowait(...,
1599/// reduceData, shuffleReduceFn, interWarpCpyFn,
1600/// scratchpadCopyFn, loadAndReduceFn)
1601///
1602/// 'scratchpadCopyFn' is a helper that stores reduced
1603/// data from the team master to a scratchpad array in
1604/// global memory.
1605///
1606/// 'loadAndReduceFn' is a helper that loads data from
1607/// the scratchpad array and reduces it with the input
1608/// operand.
1609///
1610/// These compiler generated functions hide address
1611/// calculation and alignment information from the runtime.
1612/// 5. if ret == 1:
1613/// The team master of the last team stores the reduced
1614/// result to the globals in memory.
1615/// foo += reduceData.foo; bar *= reduceData.bar
1616///
1617///
1618/// Warp Reduction Algorithms
1619///
1620/// On the warp level, we have three algorithms implemented in the
1621/// OpenMP runtime depending on the number of active lanes:
1622///
1623/// Full Warp Reduction
1624///
1625/// The reduce algorithm within a warp where all lanes are active
1626/// is implemented in the runtime as follows:
1627///
1628/// full_warp_reduce(void *reduce_data,
1629/// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
1630/// for (int offset = WARPSIZE/2; offset > 0; offset /= 2)
1631/// ShuffleReduceFn(reduce_data, 0, offset, 0);
1632/// }
1633///
1634/// The algorithm completes in log(2, WARPSIZE) steps.
1635///
1636/// 'ShuffleReduceFn' is used here with lane_id set to 0 because it is
1637/// not used therefore we save instructions by not retrieving lane_id
1638/// from the corresponding special registers. The 4th parameter, which
1639/// represents the version of the algorithm being used, is set to 0 to
1640/// signify full warp reduction.
1641///
1642/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
1643///
1644/// #reduce_elem refers to an element in the local lane's data structure
1645/// #remote_elem is retrieved from a remote lane
1646/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
1647/// reduce_elem = reduce_elem REDUCE_OP remote_elem;
1648///
1649/// Contiguous Partial Warp Reduction
1650///
1651/// This reduce algorithm is used within a warp where only the first
1652/// 'n' (n <= WARPSIZE) lanes are active. It is typically used when the
1653/// number of OpenMP threads in a parallel region is not a multiple of
1654/// WARPSIZE. The algorithm is implemented in the runtime as follows:
1655///
1656/// void
1657/// contiguous_partial_reduce(void *reduce_data,
1658/// kmp_ShuffleReductFctPtr ShuffleReduceFn,
1659/// int size, int lane_id) {
1660/// int curr_size;
1661/// int offset;
1662/// curr_size = size;
1663/// mask = curr_size/2;
1664/// while (offset>0) {
1665/// ShuffleReduceFn(reduce_data, lane_id, offset, 1);
1666/// curr_size = (curr_size+1)/2;
1667/// offset = curr_size/2;
1668/// }
1669/// }
1670///
1671/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
1672///
1673/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
1674/// if (lane_id < offset)
1675/// reduce_elem = reduce_elem REDUCE_OP remote_elem
1676/// else
1677/// reduce_elem = remote_elem
1678///
1679/// This algorithm assumes that the data to be reduced are located in a
1680/// contiguous subset of lanes starting from the first. When there is
1681/// an odd number of active lanes, the data in the last lane is not
1682/// aggregated with any other lane's dat but is instead copied over.
1683///
1684/// Dispersed Partial Warp Reduction
1685///
1686/// This algorithm is used within a warp when any discontiguous subset of
1687/// lanes are active. It is used to implement the reduction operation
1688/// across lanes in an OpenMP simd region or in a nested parallel region.
1689///
1690/// void
1691/// dispersed_partial_reduce(void *reduce_data,
1692/// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
1693/// int size, remote_id;
1694/// int logical_lane_id = number_of_active_lanes_before_me() * 2;
1695/// do {
1696/// remote_id = next_active_lane_id_right_after_me();
1697/// # the above function returns 0 of no active lane
1698/// # is present right after the current lane.
1699/// size = number_of_active_lanes_in_this_warp();
1700/// logical_lane_id /= 2;
1701/// ShuffleReduceFn(reduce_data, logical_lane_id,
1702/// remote_id-1-threadIdx.x, 2);
1703/// } while (logical_lane_id % 2 == 0 && size > 1);
1704/// }
1705///
1706/// There is no assumption made about the initial state of the reduction.
1707/// Any number of lanes (>=1) could be active at any position. The reduction
1708/// result is returned in the first active lane.
1709///
1710/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
1711///
1712/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
1713/// if (lane_id % 2 == 0 && offset > 0)
1714/// reduce_elem = reduce_elem REDUCE_OP remote_elem
1715/// else
1716/// reduce_elem = remote_elem
1717///
1718///
1719/// Intra-Team Reduction
1720///
1721/// This function, as implemented in the runtime call
1722/// '__kmpc_nvptx_parallel_reduce_nowait_v2', aggregates data across OpenMP
1723/// threads in a team. It first reduces within a warp using the
1724/// aforementioned algorithms. We then proceed to gather all such
1725/// reduced values at the first warp.
1726///
1727/// The runtime makes use of the function 'InterWarpCpyFn', which copies
1728/// data from each of the "warp master" (zeroth lane of each warp, where
1729/// warp-reduced data is held) to the zeroth warp. This step reduces (in
1730/// a mathematical sense) the problem of reduction across warp masters in
1731/// a block to the problem of warp reduction.
1732///
1733///
1734/// Inter-Team Reduction
1735///
1736/// Once a team has reduced its data to a single value, it is stored in
1737/// a global scratchpad array. Since each team has a distinct slot, this
1738/// can be done without locking.
1739///
1740/// The last team to write to the scratchpad array proceeds to reduce the
1741/// scratchpad array. One or more workers in the last team use the helper
1742/// 'loadAndReduceDataFn' to load and reduce values from the array, i.e.,
1743/// the k'th worker reduces every k'th element.
1744///
1745/// Finally, a call is made to '__kmpc_nvptx_parallel_reduce_nowait_v2' to
1746/// reduce across workers and compute a globally reduced value.
1747///
1748void CGOpenMPRuntimeGPU::emitReduction(
1749 CodeGenFunction &CGF, SourceLocation Loc, ArrayRef<const Expr *> Privates,
1750 ArrayRef<const Expr *> LHSExprs, ArrayRef<const Expr *> RHSExprs,
1751 ArrayRef<const Expr *> ReductionOps, ReductionOptionsTy Options) {
1752 if (!CGF.HaveInsertPoint())
1753 return;
1754
1755 bool ParallelReduction = isOpenMPParallelDirective(DKind: Options.ReductionKind);
1756 bool TeamsReduction = isOpenMPTeamsDirective(DKind: Options.ReductionKind);
1757
1758 if (Options.SimpleReduction) {
1759 assert(!TeamsReduction && !ParallelReduction &&
1760 "Invalid reduction selection in emitReduction.");
1761 (void)ParallelReduction;
1762 CGOpenMPRuntime::emitReduction(CGF, Loc, Privates, LHSExprs, RHSExprs,
1763 ReductionOps, Options);
1764 return;
1765 }
1766
1767 llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> VarFieldMap;
1768 llvm::SmallVector<const ValueDecl *, 4> PrivatesReductions(Privates.size());
1769 int Cnt = 0;
1770 for (const Expr *DRE : Privates) {
1771 PrivatesReductions[Cnt] = cast<DeclRefExpr>(Val: DRE)->getDecl();
1772 ++Cnt;
1773 }
1774 const RecordDecl *ReductionRec = ::buildRecordForGlobalizedVars(
1775 C&: CGM.getContext(), EscapedDecls: PrivatesReductions, EscapedDeclsForTeams: {}, MappedDeclsFields&: VarFieldMap, BufSize: 1);
1776
1777 // The atomic cross-team reduction fast path is opt-in. Hand each eligible
1778 // scalar reduction an atomic combiner; createReductionsGPU uses the atomic
1779 // path only if every reduction in the set has one. Track whether that holds
1780 // so we can skip the (then unused) per-team buffer registration.
1781 bool UseAtomicReduction =
1782 TeamsReduction && CGM.getLangOpts().OpenMPTargetAtomicReduction;
1783 bool AllAtomicable = UseAtomicReduction;
1784
1785 // Source location for the ident struct
1786 llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
1787
1788 using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
1789 InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
1790 CGF.AllocaInsertPt->getIterator());
1791 InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
1792 CGF.Builder.GetInsertPoint());
1793 llvm::OpenMPIRBuilder::LocationDescription OmpLoc(
1794 CodeGenIP, CGF.SourceLocToDebugLoc(Location: Loc));
1795 llvm::SmallVector<llvm::OpenMPIRBuilder::ReductionInfo, 2> ReductionInfos;
1796
1797 CodeGenFunction::OMPPrivateScope Scope(CGF);
1798 unsigned Idx = 0;
1799 for (const Expr *Private : Privates) {
1800 llvm::Type *ElementType;
1801 llvm::Value *Variable;
1802 llvm::Value *PrivateVariable;
1803 llvm::OpenMPIRBuilder::ReductionGenAtomicCBTy AtomicReductionGen = nullptr;
1804 ElementType = CGF.ConvertTypeForMem(T: Private->getType());
1805 const auto *RHSVar =
1806 cast<VarDecl>(Val: cast<DeclRefExpr>(Val: RHSExprs[Idx])->getDecl());
1807 PrivateVariable = CGF.GetAddrOfLocalVar(VD: RHSVar).emitRawPointer(CGF);
1808 const auto *LHSVar =
1809 cast<VarDecl>(Val: cast<DeclRefExpr>(Val: LHSExprs[Idx])->getDecl());
1810 Variable = CGF.GetAddrOfLocalVar(VD: LHSVar).emitRawPointer(CGF);
1811 llvm::OpenMPIRBuilder::EvalKind EvalKind;
1812 switch (CGF.getEvaluationKind(T: Private->getType())) {
1813 case TEK_Scalar:
1814 EvalKind = llvm::OpenMPIRBuilder::EvalKind::Scalar;
1815 break;
1816 case TEK_Complex:
1817 EvalKind = llvm::OpenMPIRBuilder::EvalKind::Complex;
1818 break;
1819 case TEK_Aggregate:
1820 EvalKind = llvm::OpenMPIRBuilder::EvalKind::Aggregate;
1821 break;
1822 }
1823 auto ReductionGen = [&](InsertPointTy CodeGenIP, unsigned I,
1824 llvm::Value **LHSPtr, llvm::Value **RHSPtr,
1825 llvm::Function *NewFunc) {
1826 CGF.Builder.restoreIP(IP: CodeGenIP);
1827 auto *CurFn = CGF.CurFn;
1828 CGF.CurFn = NewFunc;
1829
1830 // The helper has no DISubprogram of its own, so a debug location here
1831 // would name the enclosing function's scope, which is invalid IR.
1832 // Suppress them, as the other OpenMPIRBuilder-generated helpers do.
1833 llvm::DebugLoc SavedDebugLoc = CGF.Builder.getCurrentDebugLocation();
1834 CGF.Builder.SetCurrentDebugLocation(llvm::DebugLoc());
1835 CGF.disableDebugInfo();
1836
1837 *LHSPtr = CGF.GetAddrOfLocalVar(
1838 VD: cast<VarDecl>(Val: cast<DeclRefExpr>(Val: LHSExprs[I])->getDecl()))
1839 .emitRawPointer(CGF);
1840 *RHSPtr = CGF.GetAddrOfLocalVar(
1841 VD: cast<VarDecl>(Val: cast<DeclRefExpr>(Val: RHSExprs[I])->getDecl()))
1842 .emitRawPointer(CGF);
1843
1844 emitSingleReductionCombiner(CGF, ReductionOp: ReductionOps[I], PrivateRef: Privates[I],
1845 LHS: cast<DeclRefExpr>(Val: LHSExprs[I]),
1846 RHS: cast<DeclRefExpr>(Val: RHSExprs[I]));
1847
1848 CGF.enableDebugInfo();
1849 CGF.Builder.SetCurrentDebugLocation(SavedDebugLoc);
1850 CGF.CurFn = CurFn;
1851
1852 return InsertPointTy(CGF.Builder.GetInsertBlock(),
1853 CGF.Builder.GetInsertPoint());
1854 };
1855
1856 // For the atomic fast path, hand this reduction an atomic combiner if it is
1857 // a scalar with a direct atomicrmw; otherwise the set is not fully
1858 // atomicable and falls back to the buffer path.
1859 if (UseAtomicReduction) {
1860 std::optional<llvm::AtomicRMWInst::BinOp> AtomicOp;
1861 if (EvalKind == llvm::OpenMPIRBuilder::EvalKind::Scalar) {
1862 if (std::optional<BinaryOperatorKind> BOK =
1863 getReductionBinOpKind(ReductionOp: ReductionOps[Idx]))
1864 AtomicOp = getReductionAtomicRMWOp(BOK: *BOK, Ty: Private->getType());
1865 }
1866 if (!AtomicOp) {
1867 AllAtomicable = false;
1868 } else {
1869 llvm::AtomicRMWInst::BinOp Op = *AtomicOp;
1870 llvm::Align Alignment =
1871 CGM.getModule().getDataLayout().getPrefTypeAlign(Ty: ElementType);
1872 // Device (agent) scope suffices: all teams accumulate on-device and the
1873 // host reads the result only after the kernel (via map-back), so the
1874 // far costlier system scope is unnecessary. The
1875 // no.fine.grained/no.remote memory metadata is omitted so the atomic
1876 // stays correct under USM.
1877 llvm::SyncScope::ID SSID = CGF.getTargetHooks().getLLVMSyncScopeID(
1878 LangOpts: CGF.getLangOpts(), Scope: SyncScope::DeviceScope,
1879 Ordering: llvm::AtomicOrdering::Monotonic, Ctx&: CGF.getLLVMContext());
1880 AtomicReductionGen = [Op, Alignment,
1881 SSID](InsertPointTy IP, llvm::Type *EltTy,
1882 llvm::Value *LHS, llvm::Value *RHS)
1883 -> llvm::OpenMPIRBuilder::InsertPointOrErrorTy {
1884 llvm::IRBuilder<> Builder(IP.getBlock(), IP.getPoint());
1885 llvm::Value *Val = Builder.CreateLoad(Ty: EltTy, Ptr: RHS);
1886 Builder.CreateAtomicRMW(Op, Ptr: LHS, Val, Align: Alignment,
1887 Ordering: llvm::AtomicOrdering::Monotonic, SSID);
1888 return InsertPointTy(Builder.GetInsertBlock(),
1889 Builder.GetInsertPoint());
1890 };
1891 }
1892 }
1893
1894 ReductionInfos.emplace_back(Args: llvm::OpenMPIRBuilder::ReductionInfo(
1895 ElementType, Variable, PrivateVariable, EvalKind,
1896 /*ReductionGen=*/nullptr, ReductionGen, AtomicReductionGen,
1897 /*DataPtrPtrGen=*/nullptr));
1898 Idx++;
1899 }
1900
1901 // The atomic path folds directly into the mapped variable and needs no
1902 // per-team buffer; register the record for buffer allocation otherwise.
1903 if (TeamsReduction && !AllAtomicable)
1904 TeamsReductions.push_back(Elt: ReductionRec);
1905
1906 bool IsSPMD = getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD;
1907 llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
1908 cantFail(ValOrErr: OMPBuilder.createReductionsGPU(
1909 Loc: OmpLoc, AllocaIP, CodeGenIP, ReductionInfos, /*IsByRef=*/{}, IsNoWait: false,
1910 IsTeamsReduction: TeamsReduction, IsSPMD,
1911 ReductionGenCBKind: llvm::OpenMPIRBuilder::ReductionGenCBKind::Clang,
1912 GridValue: CGF.getTarget().getGridValue(), SrcLocInfo: RTLoc));
1913 CGF.Builder.restoreIP(IP: AfterIP);
1914}
1915
1916const VarDecl *
1917CGOpenMPRuntimeGPU::translateParameter(const FieldDecl *FD,
1918 const VarDecl *NativeParam) const {
1919 if (!NativeParam->getType()->isReferenceType())
1920 return NativeParam;
1921 QualType ArgType = NativeParam->getType();
1922 QualifierCollector QC;
1923 const Type *NonQualTy = QC.strip(type: ArgType);
1924 QualType PointeeTy = cast<ReferenceType>(Val: NonQualTy)->getPointeeType();
1925 if (const auto *Attr = FD->getAttr<OMPCaptureKindAttr>()) {
1926 if (Attr->getCaptureKind() == OMPC_map) {
1927 PointeeTy = CGM.getContext().getAddrSpaceQualType(T: PointeeTy,
1928 AddressSpace: LangAS::opencl_global);
1929 }
1930 }
1931 ArgType = CGM.getContext().getPointerType(T: PointeeTy);
1932 QC.addRestrict();
1933 ArgType = QC.apply(Context: CGM.getContext(), QT: ArgType);
1934 if (isa<ImplicitParamDecl>(Val: NativeParam))
1935 return ImplicitParamDecl::Create(
1936 C&: CGM.getContext(), /*DC=*/nullptr, IdLoc: NativeParam->getLocation(),
1937 Id: NativeParam->getIdentifier(), T: ArgType, ParamKind: ImplicitParamKind::Other);
1938 return ParmVarDecl::Create(
1939 C&: CGM.getContext(),
1940 DC: const_cast<DeclContext *>(NativeParam->getDeclContext()),
1941 StartLoc: NativeParam->getBeginLoc(), IdLoc: NativeParam->getLocation(),
1942 Id: NativeParam->getIdentifier(), T: ArgType,
1943 /*TInfo=*/nullptr, S: SC_None, /*DefArg=*/nullptr);
1944}
1945
1946Address
1947CGOpenMPRuntimeGPU::getParameterAddress(CodeGenFunction &CGF,
1948 const VarDecl *NativeParam,
1949 const VarDecl *TargetParam) const {
1950 assert(NativeParam != TargetParam &&
1951 NativeParam->getType()->isReferenceType() &&
1952 "Native arg must not be the same as target arg.");
1953 Address LocalAddr = CGF.GetAddrOfLocalVar(VD: TargetParam);
1954 QualType NativeParamType = NativeParam->getType();
1955 QualifierCollector QC;
1956 const Type *NonQualTy = QC.strip(type: NativeParamType);
1957 QualType NativePointeeTy = cast<ReferenceType>(Val: NonQualTy)->getPointeeType();
1958 unsigned NativePointeeAddrSpace =
1959 CGF.getTypes().getTargetAddressSpace(T: NativePointeeTy);
1960 QualType TargetTy = TargetParam->getType();
1961 llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(Addr: LocalAddr, /*Volatile=*/false,
1962 Ty: TargetTy, Loc: SourceLocation());
1963 // Cast to native address space.
1964 TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
1965 V: TargetAddr,
1966 DestTy: llvm::PointerType::get(C&: CGF.getLLVMContext(), AddressSpace: NativePointeeAddrSpace));
1967 Address NativeParamAddr = CGF.CreateMemTemp(T: NativeParamType);
1968 CGF.EmitStoreOfScalar(Value: TargetAddr, Addr: NativeParamAddr, /*Volatile=*/false,
1969 Ty: NativeParamType);
1970 return NativeParamAddr;
1971}
1972
1973void CGOpenMPRuntimeGPU::emitOutlinedFunctionCall(
1974 CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn,
1975 ArrayRef<llvm::Value *> Args) const {
1976 SmallVector<llvm::Value *, 4> TargetArgs;
1977 TargetArgs.reserve(N: Args.size());
1978 auto *FnType = OutlinedFn.getFunctionType();
1979 for (unsigned I = 0, E = Args.size(); I < E; ++I) {
1980 if (FnType->isVarArg() && FnType->getNumParams() <= I) {
1981 TargetArgs.append(in_start: std::next(x: Args.begin(), n: I), in_end: Args.end());
1982 break;
1983 }
1984 llvm::Type *TargetType = FnType->getParamType(i: I);
1985 llvm::Value *NativeArg = Args[I];
1986 if (!TargetType->isPointerTy()) {
1987 TargetArgs.emplace_back(Args&: NativeArg);
1988 continue;
1989 }
1990 TargetArgs.emplace_back(
1991 Args: CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(V: NativeArg, DestTy: TargetType));
1992 }
1993 CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, Args: TargetArgs);
1994}
1995
1996/// Emit function which wraps the outline parallel region
1997/// and controls the arguments which are passed to this function.
1998/// The wrapper ensures that the outlined function is called
1999/// with the correct arguments when data is shared.
2000llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
2001 llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) {
2002 ASTContext &Ctx = CGM.getContext();
2003 const auto &CS = *D.getCapturedStmt(RegionKind: OMPD_parallel);
2004
2005 // Create a function that takes as argument the source thread.
2006 FunctionArgList WrapperArgs;
2007 QualType Int16QTy =
2008 Ctx.getIntTypeForBitwidth(/*DestWidth=*/16, /*Signed=*/false);
2009 QualType Int32QTy =
2010 Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/false);
2011 auto *ParallelLevelArg = ImplicitParamDecl::Create(
2012 C&: Ctx, /*DC=*/nullptr, IdLoc: D.getBeginLoc(),
2013 /*Id=*/nullptr, T: Int16QTy, ParamKind: ImplicitParamKind::Other);
2014 auto *WrapperArg = ImplicitParamDecl::Create(
2015 C&: Ctx, /*DC=*/nullptr, IdLoc: D.getBeginLoc(),
2016 /*Id=*/nullptr, T: Int32QTy, ParamKind: ImplicitParamKind::Other);
2017 WrapperArgs.emplace_back(Args&: ParallelLevelArg);
2018 WrapperArgs.emplace_back(Args&: WrapperArg);
2019
2020 const CGFunctionInfo &CGFI =
2021 CGM.getTypes().arrangeBuiltinFunctionDeclaration(resultType: Ctx.VoidTy, args: WrapperArgs);
2022
2023 auto *Fn = llvm::Function::Create(
2024 Ty: CGM.getTypes().GetFunctionType(Info: CGFI), Linkage: llvm::GlobalValue::InternalLinkage,
2025 N: Twine(OutlinedParallelFn->getName(), "_wrapper"), M: &CGM.getModule());
2026
2027 // Ensure we do not inline the function. This is trivially true for the ones
2028 // passed to __kmpc_fork_call but the ones calles in serialized regions
2029 // could be inlined. This is not a perfect but it is closer to the invariant
2030 // we want, namely, every data environment starts with a new function.
2031 // TODO: We should pass the if condition to the runtime function and do the
2032 // handling there. Much cleaner code.
2033 Fn->addFnAttr(Kind: llvm::Attribute::NoInline);
2034
2035 CGM.SetInternalFunctionAttributes(GD: GlobalDecl(), F: Fn, FI: CGFI);
2036 Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
2037 Fn->setDoesNotRecurse();
2038
2039 CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
2040 CGF.StartFunction(GD: GlobalDecl(), RetTy: Ctx.VoidTy, Fn, FnInfo: CGFI, Args: WrapperArgs,
2041 Loc: D.getBeginLoc(), StartLoc: D.getBeginLoc());
2042
2043 const auto *RD = CS.getCapturedRecordDecl();
2044 auto CurField = RD->field_begin();
2045
2046 Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(Ty: CGF.Int32Ty,
2047 /*Name=*/".zero.addr");
2048 CGF.Builder.CreateStore(Val: CGF.Builder.getInt32(/*C*/ 0), Addr: ZeroAddr);
2049 // Get the array of arguments.
2050 SmallVector<llvm::Value *, 8> Args;
2051
2052 Args.emplace_back(Args: CGF.GetAddrOfLocalVar(VD: WrapperArg).emitRawPointer(CGF));
2053 Args.emplace_back(Args: ZeroAddr.emitRawPointer(CGF));
2054
2055 CGBuilderTy &Bld = CGF.Builder;
2056 auto CI = CS.capture_begin();
2057
2058 // Use global memory for data sharing.
2059 // Handle passing of global args to workers.
2060 RawAddress GlobalArgs =
2061 CGF.CreateDefaultAlignTempAlloca(Ty: CGF.VoidPtrPtrTy, Name: "global_args");
2062 llvm::Value *GlobalArgsPtr = GlobalArgs.getPointer();
2063 llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};
2064 CGF.EmitRuntimeCall(callee: OMPBuilder.getOrCreateRuntimeFunction(
2065 M&: CGM.getModule(), FnID: OMPRTL___kmpc_get_shared_variables),
2066 args: DataSharingArgs);
2067
2068 // Retrieve the shared variables from the list of references returned
2069 // by the runtime. Pass the variables to the outlined function.
2070 Address SharedArgListAddress = Address::invalid();
2071 if (CS.capture_size() > 0 ||
2072 isOpenMPLoopBoundSharingDirective(Kind: D.getDirectiveKind())) {
2073 SharedArgListAddress = CGF.EmitLoadOfPointer(
2074 Ptr: GlobalArgs, PtrTy: CGF.getContext()
2075 .getPointerType(T: CGF.getContext().VoidPtrTy)
2076 .castAs<PointerType>());
2077 }
2078 unsigned Idx = 0;
2079 if (isOpenMPLoopBoundSharingDirective(Kind: D.getDirectiveKind())) {
2080 Address Src = Bld.CreateConstInBoundsGEP(Addr: SharedArgListAddress, Index: Idx);
2081 Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
2082 Addr: Src, Ty: Bld.getPtrTy(AddrSpace: 0), ElementTy: CGF.SizeTy);
2083 llvm::Value *LB = CGF.EmitLoadOfScalar(
2084 Addr: TypedAddress,
2085 /*Volatile=*/false,
2086 Ty: CGF.getContext().getPointerType(T: CGF.getContext().getSizeType()),
2087 Loc: cast<OMPLoopDirective>(Val: D).getLowerBoundVariable()->getExprLoc());
2088 Args.emplace_back(Args&: LB);
2089 ++Idx;
2090 Src = Bld.CreateConstInBoundsGEP(Addr: SharedArgListAddress, Index: Idx);
2091 TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(Addr: Src, Ty: Bld.getPtrTy(AddrSpace: 0),
2092 ElementTy: CGF.SizeTy);
2093 llvm::Value *UB = CGF.EmitLoadOfScalar(
2094 Addr: TypedAddress,
2095 /*Volatile=*/false,
2096 Ty: CGF.getContext().getPointerType(T: CGF.getContext().getSizeType()),
2097 Loc: cast<OMPLoopDirective>(Val: D).getUpperBoundVariable()->getExprLoc());
2098 Args.emplace_back(Args&: UB);
2099 ++Idx;
2100 }
2101 if (CS.capture_size() > 0) {
2102 ASTContext &CGFContext = CGF.getContext();
2103 for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CI, ++CurField) {
2104 QualType ElemTy = CurField->getType();
2105 Address Src = Bld.CreateConstInBoundsGEP(Addr: SharedArgListAddress, Index: I + Idx);
2106 Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
2107 Addr: Src, Ty: CGF.ConvertTypeForMem(T: CGFContext.getPointerType(T: ElemTy)),
2108 ElementTy: CGF.ConvertTypeForMem(T: ElemTy));
2109 llvm::Value *Arg = CGF.EmitLoadOfScalar(Addr: TypedAddress,
2110 /*Volatile=*/false,
2111 Ty: CGFContext.getPointerType(T: ElemTy),
2112 Loc: CI->getLocation());
2113 if (CI->capturesVariableByCopy() &&
2114 !CI->getCapturedVar()->getType()->isAnyPointerType()) {
2115 Arg = castValueToType(CGF, Val: Arg, ValTy: ElemTy, CastTy: CGFContext.getUIntPtrType(),
2116 Loc: CI->getLocation());
2117 }
2118 Args.emplace_back(Args&: Arg);
2119 }
2120 }
2121
2122 emitOutlinedFunctionCall(CGF, Loc: D.getBeginLoc(), OutlinedFn: OutlinedParallelFn, Args);
2123 CGF.FinishFunction();
2124 return Fn;
2125}
2126
2127void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,
2128 const Decl *D) {
2129 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
2130 return;
2131
2132 assert(D && "Expected function or captured|block decl.");
2133 assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&
2134 "Function is registered already.");
2135 assert((!TeamAndReductions.first || TeamAndReductions.first == D) &&
2136 "Team is set but not processed.");
2137 const Stmt *Body = nullptr;
2138 bool NeedToDelayGlobalization = false;
2139 if (const auto *FD = dyn_cast<FunctionDecl>(Val: D)) {
2140 Body = FD->getBody();
2141 } else if (const auto *BD = dyn_cast<BlockDecl>(Val: D)) {
2142 Body = BD->getBody();
2143 } else if (const auto *CD = dyn_cast<CapturedDecl>(Val: D)) {
2144 Body = CD->getBody();
2145 NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;
2146 if (NeedToDelayGlobalization &&
2147 getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
2148 return;
2149 }
2150 if (!Body)
2151 return;
2152 CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);
2153 VarChecker.Visit(S: Body);
2154 const RecordDecl *GlobalizedVarsRecord =
2155 VarChecker.getGlobalizedRecord(IsInTTDRegion);
2156 TeamAndReductions.first = nullptr;
2157 TeamAndReductions.second.clear();
2158 ArrayRef<const ValueDecl *> EscapedVariableLengthDecls =
2159 VarChecker.getEscapedVariableLengthDecls();
2160 ArrayRef<const ValueDecl *> DelayedVariableLengthDecls =
2161 VarChecker.getDelayedVariableLengthDecls();
2162 if (!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty() &&
2163 DelayedVariableLengthDecls.empty())
2164 return;
2165 auto I = FunctionGlobalizedDecls.try_emplace(Key: CGF.CurFn).first;
2166 I->getSecond().MappedParams =
2167 std::make_unique<CodeGenFunction::OMPMapVars>();
2168 I->getSecond().EscapedParameters.insert(
2169 I: VarChecker.getEscapedParameters().begin(),
2170 E: VarChecker.getEscapedParameters().end());
2171 I->getSecond().EscapedVariableLengthDecls.append(
2172 in_start: EscapedVariableLengthDecls.begin(), in_end: EscapedVariableLengthDecls.end());
2173 I->getSecond().DelayedVariableLengthDecls.append(
2174 in_start: DelayedVariableLengthDecls.begin(), in_end: DelayedVariableLengthDecls.end());
2175 DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
2176 for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
2177 assert(VD->isCanonicalDecl() && "Expected canonical declaration");
2178 Data.try_emplace(Key: VD);
2179 }
2180 if (!NeedToDelayGlobalization) {
2181 emitGenericVarsProlog(CGF, Loc: D->getBeginLoc());
2182 struct GlobalizationScope final : EHScopeStack::Cleanup {
2183 GlobalizationScope() = default;
2184
2185 void Emit(CodeGenFunction &CGF, Flags flags) override {
2186 static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
2187 .emitGenericVarsEpilog(CGF);
2188 }
2189 };
2190 CGF.EHStack.pushCleanup<GlobalizationScope>(Kind: NormalAndEHCleanup);
2191 }
2192}
2193
2194Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
2195 const VarDecl *VD) {
2196 if (VD && VD->hasAttr<OMPAllocateDeclAttr>()) {
2197 const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
2198 auto AS = LangAS::Default;
2199 switch (A->getAllocatorType()) {
2200 case OMPAllocateDeclAttr::OMPNullMemAlloc:
2201 case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
2202 case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
2203 case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
2204 break;
2205 case OMPAllocateDeclAttr::OMPThreadMemAlloc:
2206 return Address::invalid();
2207 case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
2208 // TODO: implement aupport for user-defined allocators.
2209 return Address::invalid();
2210 case OMPAllocateDeclAttr::OMPConstMemAlloc:
2211 AS = LangAS::cuda_constant;
2212 break;
2213 case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
2214 AS = LangAS::cuda_shared;
2215 break;
2216 case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
2217 case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
2218 break;
2219 }
2220 llvm::Type *VarTy = CGF.ConvertTypeForMem(T: VD->getType());
2221 auto *GV = new llvm::GlobalVariable(
2222 CGM.getModule(), VarTy, /*isConstant=*/false,
2223 llvm::GlobalValue::InternalLinkage, llvm::PoisonValue::get(T: VarTy),
2224 VD->getName(),
2225 /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
2226 CGM.getContext().getTargetAddressSpace(AS));
2227 CharUnits Align = CGM.getContext().getDeclAlign(D: VD);
2228 GV->setAlignment(Align.getAsAlign());
2229 return Address(
2230 CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
2231 V: GV, DestTy: CGF.Builder.getPtrTy(AddrSpace: CGM.getContext().getTargetAddressSpace(
2232 AS: VD->getType().getAddressSpace()))),
2233 VarTy, Align);
2234 }
2235
2236 if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
2237 return Address::invalid();
2238
2239 VD = VD->getCanonicalDecl();
2240 auto I = FunctionGlobalizedDecls.find(Val: CGF.CurFn);
2241 if (I == FunctionGlobalizedDecls.end())
2242 return Address::invalid();
2243 auto VDI = I->getSecond().LocalVarData.find(Key: VD);
2244 if (VDI != I->getSecond().LocalVarData.end())
2245 return VDI->second.PrivateAddr;
2246 if (VD->hasAttrs()) {
2247 for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),
2248 E(VD->attr_end());
2249 IT != E; ++IT) {
2250 auto VDI = I->getSecond().LocalVarData.find(
2251 Key: cast<VarDecl>(Val: cast<DeclRefExpr>(Val: IT->getRef())->getDecl())
2252 ->getCanonicalDecl());
2253 if (VDI != I->getSecond().LocalVarData.end())
2254 return VDI->second.PrivateAddr;
2255 }
2256 }
2257
2258 return Address::invalid();
2259}
2260
2261void CGOpenMPRuntimeGPU::functionFinished(CodeGenFunction &CGF) {
2262 FunctionGlobalizedDecls.erase(Val: CGF.CurFn);
2263 CGOpenMPRuntime::functionFinished(CGF);
2264}
2265
2266void CGOpenMPRuntimeGPU::getDefaultDistScheduleAndChunk(
2267 CodeGenFunction &CGF, const OMPLoopDirective &S,
2268 OpenMPDistScheduleClauseKind &ScheduleKind,
2269 llvm::Value *&Chunk) const {
2270 auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
2271 if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
2272 ScheduleKind = OMPC_DIST_SCHEDULE_static;
2273 Chunk = CGF.EmitScalarConversion(
2274 Src: RT.getGPUNumThreads(CGF),
2275 SrcTy: CGF.getContext().getIntTypeForBitwidth(DestWidth: 32, /*Signed=*/0),
2276 DstTy: S.getIterationVariable()->getType(), Loc: S.getBeginLoc());
2277 return;
2278 }
2279 CGOpenMPRuntime::getDefaultDistScheduleAndChunk(
2280 CGF, S, ScheduleKind, Chunk);
2281}
2282
2283void CGOpenMPRuntimeGPU::getDefaultScheduleAndChunk(
2284 CodeGenFunction &CGF, const OMPLoopDirective &S,
2285 OpenMPScheduleClauseKind &ScheduleKind,
2286 const Expr *&ChunkExpr) const {
2287 ScheduleKind = OMPC_SCHEDULE_static;
2288 // Chunk size is 1 in this case.
2289 llvm::APInt ChunkSize(32, 1);
2290 ChunkExpr = IntegerLiteral::Create(C: CGF.getContext(), V: ChunkSize,
2291 type: CGF.getContext().getIntTypeForBitwidth(DestWidth: 32, /*Signed=*/0),
2292 l: SourceLocation());
2293}
2294
2295void CGOpenMPRuntimeGPU::adjustTargetSpecificDataForLambdas(
2296 CodeGenFunction &CGF, const OMPExecutableDirective &D) const {
2297 assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
2298 " Expected target-based directive.");
2299 const CapturedStmt *CS = D.getCapturedStmt(RegionKind: OMPD_target);
2300 for (const CapturedStmt::Capture &C : CS->captures()) {
2301 // Capture variables captured by reference in lambdas for target-based
2302 // directives.
2303 if (!C.capturesVariable())
2304 continue;
2305 const VarDecl *VD = C.getCapturedVar();
2306 const auto *RD = VD->getType()
2307 .getCanonicalType()
2308 .getNonReferenceType()
2309 ->getAsCXXRecordDecl();
2310 if (!RD || !RD->isLambda())
2311 continue;
2312 Address VDAddr = CGF.GetAddrOfLocalVar(VD);
2313 LValue VDLVal;
2314 if (VD->getType().getCanonicalType()->isReferenceType())
2315 VDLVal = CGF.EmitLoadOfReferenceLValue(RefAddr: VDAddr, RefTy: VD->getType());
2316 else
2317 VDLVal = CGF.MakeAddrLValue(
2318 Addr: VDAddr, T: VD->getType().getCanonicalType().getNonReferenceType());
2319 llvm::DenseMap<const ValueDecl *, FieldDecl *> Captures;
2320 FieldDecl *ThisCapture = nullptr;
2321 RD->getCaptureFields(Captures, ThisCapture);
2322 if (ThisCapture && CGF.CapturedStmtInfo->isCXXThisExprCaptured()) {
2323 LValue ThisLVal =
2324 CGF.EmitLValueForFieldInitialization(Base: VDLVal, Field: ThisCapture);
2325 llvm::Value *CXXThis = CGF.LoadCXXThis();
2326 CGF.EmitStoreOfScalar(value: CXXThis, lvalue: ThisLVal);
2327 }
2328 for (const LambdaCapture &LC : RD->captures()) {
2329 if (LC.getCaptureKind() != LCK_ByRef)
2330 continue;
2331 const ValueDecl *VD = LC.getCapturedVar();
2332 // FIXME: For now VD is always a VarDecl because OpenMP does not support
2333 // capturing structured bindings in lambdas yet.
2334 if (!CS->capturesVariable(Var: cast<VarDecl>(Val: VD)))
2335 continue;
2336 auto It = Captures.find(Val: VD);
2337 assert(It != Captures.end() && "Found lambda capture without field.");
2338 LValue VarLVal = CGF.EmitLValueForFieldInitialization(Base: VDLVal, Field: It->second);
2339 Address VDAddr = CGF.GetAddrOfLocalVar(VD: cast<VarDecl>(Val: VD));
2340 if (VD->getType().getCanonicalType()->isReferenceType())
2341 VDAddr = CGF.EmitLoadOfReferenceLValue(RefAddr: VDAddr,
2342 RefTy: VD->getType().getCanonicalType())
2343 .getAddress();
2344 CGF.EmitStoreOfScalar(value: VDAddr.emitRawPointer(CGF), lvalue: VarLVal);
2345 }
2346 }
2347}
2348
2349bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
2350 LangAS &AS) {
2351 if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
2352 return false;
2353 const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
2354 switch(A->getAllocatorType()) {
2355 case OMPAllocateDeclAttr::OMPNullMemAlloc:
2356 case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
2357 // Not supported, fallback to the default mem space.
2358 case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
2359 case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
2360 case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
2361 case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
2362 case OMPAllocateDeclAttr::OMPThreadMemAlloc:
2363 AS = LangAS::Default;
2364 return true;
2365 case OMPAllocateDeclAttr::OMPConstMemAlloc:
2366 AS = LangAS::cuda_constant;
2367 return true;
2368 case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
2369 AS = LangAS::cuda_shared;
2370 return true;
2371 case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
2372 llvm_unreachable("Expected predefined allocator for the variables with the "
2373 "static storage.");
2374 }
2375 return false;
2376}
2377
2378/// Check to see if target architecture supports unified addressing which is
2379/// a restriction for OpenMP requires clause "unified_shared_memory".
2380void CGOpenMPRuntimeGPU::processRequiresDirective(const OMPRequiresDecl *D) {
2381 StringRef CPU = CGM.getTarget().getTargetOpts().CPU;
2382 if (CGM.getTarget().getTriple().isNVPTX() &&
2383 !llvm::NVPTX::supportsUnifiedAddressing(Kind: llvm::NVPTX::parseArch(CPU))) {
2384 for (const OMPClause *Clause : D->clauselists()) {
2385 if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
2386 CGM.getDiags().Report(Loc: Clause->getBeginLoc(),
2387 DiagID: diag::err_omp_unified_shared_memory_unsupported)
2388 << CPU;
2389 return;
2390 }
2391 }
2392 }
2393
2394 CGOpenMPRuntime::processRequiresDirective(D);
2395}
2396
2397llvm::Value *CGOpenMPRuntimeGPU::getGPUNumThreads(CodeGenFunction &CGF) {
2398 CGBuilderTy &Bld = CGF.Builder;
2399 llvm::Module *M = &CGF.CGM.getModule();
2400 const char *LocSize = "__kmpc_get_hardware_num_threads_in_block";
2401 llvm::Function *F = M->getFunction(Name: LocSize);
2402 if (!F) {
2403 F = llvm::Function::Create(Ty: llvm::FunctionType::get(Result: CGF.Int32Ty, Params: {}, isVarArg: false),
2404 Linkage: llvm::GlobalVariable::ExternalLinkage, N: LocSize,
2405 M: &CGF.CGM.getModule());
2406 }
2407 return Bld.CreateCall(Callee: F, Args: {}, Name: "nvptx_num_threads");
2408}
2409
2410llvm::Value *CGOpenMPRuntimeGPU::getGPUThreadID(CodeGenFunction &CGF) {
2411 ArrayRef<llvm::Value *> Args{};
2412 return CGF.EmitRuntimeCall(
2413 callee: OMPBuilder.getOrCreateRuntimeFunction(
2414 M&: CGM.getModule(), FnID: OMPRTL___kmpc_get_hardware_thread_id_in_block),
2415 args: Args);
2416}
2417