Bug Summary

File:tools/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
Warning:line 4787, column 8
Called C++ object pointer is null

Annotated Source Code

Press '?' to see keyboard shortcuts

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