Coverage Report

Created: 2021-08-24 07:12

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