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