Bug Summary

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