Coverage Report

Created: 2022-01-22 13:19

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