Coverage Report

Created: 2019-07-24 05:18

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