Coverage Report

Created: 2020-02-25 14:32

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