1 //===---- CGOpenMPRuntimeNVPTX.cpp - Interface to OpenMP NVPTX Runtimes ---===//
3 // The LLVM Compiler Infrastructure
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
8 //===----------------------------------------------------------------------===//
10 // This provides a class for OpenMP runtime code generation specialized to NVPTX
13 //===----------------------------------------------------------------------===//
15 #include "CGOpenMPRuntimeNVPTX.h"
16 #include "clang/AST/DeclOpenMP.h"
17 #include "CodeGenFunction.h"
18 #include "clang/AST/StmtOpenMP.h"
20 using namespace clang;
21 using namespace CodeGen;
24 enum OpenMPRTLFunctionNVPTX {
25 /// \brief Call to void __kmpc_kernel_init(kmp_int32 thread_limit);
26 OMPRTL_NVPTX__kmpc_kernel_init,
27 /// \brief Call to void __kmpc_kernel_deinit();
28 OMPRTL_NVPTX__kmpc_kernel_deinit,
29 /// \brief Call to void __kmpc_spmd_kernel_init(kmp_int32 thread_limit,
30 /// short RequiresOMPRuntime, short RequiresDataSharing);
31 OMPRTL_NVPTX__kmpc_spmd_kernel_init,
32 /// \brief Call to void __kmpc_spmd_kernel_deinit();
33 OMPRTL_NVPTX__kmpc_spmd_kernel_deinit,
34 /// \brief Call to void __kmpc_kernel_prepare_parallel(void
35 /// *outlined_function);
36 OMPRTL_NVPTX__kmpc_kernel_prepare_parallel,
37 /// \brief Call to bool __kmpc_kernel_parallel(void **outlined_function);
38 OMPRTL_NVPTX__kmpc_kernel_parallel,
39 /// \brief Call to void __kmpc_kernel_end_parallel();
40 OMPRTL_NVPTX__kmpc_kernel_end_parallel,
41 /// Call to void __kmpc_serialized_parallel(ident_t *loc, kmp_int32
43 OMPRTL_NVPTX__kmpc_serialized_parallel,
44 /// Call to void __kmpc_end_serialized_parallel(ident_t *loc, kmp_int32
46 OMPRTL_NVPTX__kmpc_end_serialized_parallel,
47 /// \brief Call to int32_t __kmpc_shuffle_int32(int32_t element,
48 /// int16_t lane_offset, int16_t warp_size);
49 OMPRTL_NVPTX__kmpc_shuffle_int32,
50 /// \brief Call to int64_t __kmpc_shuffle_int64(int64_t element,
51 /// int16_t lane_offset, int16_t warp_size);
52 OMPRTL_NVPTX__kmpc_shuffle_int64,
53 /// \brief Call to __kmpc_nvptx_parallel_reduce_nowait(kmp_int32
54 /// global_tid, kmp_int32 num_vars, size_t reduce_size, void* reduce_data,
55 /// void (*kmp_ShuffleReductFctPtr)(void *rhsData, int16_t lane_id, int16_t
56 /// lane_offset, int16_t shortCircuit),
57 /// void (*kmp_InterWarpCopyFctPtr)(void* src, int32_t warp_num));
58 OMPRTL_NVPTX__kmpc_parallel_reduce_nowait,
59 /// \brief Call to __kmpc_nvptx_teams_reduce_nowait(int32_t global_tid,
60 /// int32_t num_vars, size_t reduce_size, void *reduce_data,
61 /// void (*kmp_ShuffleReductFctPtr)(void *rhs, int16_t lane_id, int16_t
62 /// lane_offset, int16_t shortCircuit),
63 /// void (*kmp_InterWarpCopyFctPtr)(void* src, int32_t warp_num),
64 /// void (*kmp_CopyToScratchpadFctPtr)(void *reduce_data, void * scratchpad,
65 /// int32_t index, int32_t width),
66 /// void (*kmp_LoadReduceFctPtr)(void *reduce_data, void * scratchpad, int32_t
67 /// index, int32_t width, int32_t reduce))
68 OMPRTL_NVPTX__kmpc_teams_reduce_nowait,
69 /// \brief Call to __kmpc_nvptx_end_reduce_nowait(int32_t global_tid);
70 OMPRTL_NVPTX__kmpc_end_reduce_nowait
73 /// Pre(post)-action for different OpenMP constructs specialized for NVPTX.
74 class NVPTXActionTy final : public PrePostActionTy {
75 llvm::Value *EnterCallee;
76 ArrayRef<llvm::Value *> EnterArgs;
77 llvm::Value *ExitCallee;
78 ArrayRef<llvm::Value *> ExitArgs;
80 llvm::BasicBlock *ContBlock = nullptr;
83 NVPTXActionTy(llvm::Value *EnterCallee, ArrayRef<llvm::Value *> EnterArgs,
84 llvm::Value *ExitCallee, ArrayRef<llvm::Value *> ExitArgs,
85 bool Conditional = false)
86 : EnterCallee(EnterCallee), EnterArgs(EnterArgs), ExitCallee(ExitCallee),
87 ExitArgs(ExitArgs), Conditional(Conditional) {}
88 void Enter(CodeGenFunction &CGF) override {
89 llvm::Value *EnterRes = CGF.EmitRuntimeCall(EnterCallee, EnterArgs);
91 llvm::Value *CallBool = CGF.Builder.CreateIsNotNull(EnterRes);
92 auto *ThenBlock = CGF.createBasicBlock("omp_if.then");
93 ContBlock = CGF.createBasicBlock("omp_if.end");
94 // Generate the branch (If-stmt)
95 CGF.Builder.CreateCondBr(CallBool, ThenBlock, ContBlock);
96 CGF.EmitBlock(ThenBlock);
99 void Done(CodeGenFunction &CGF) {
100 // Emit the rest of blocks/branches
101 CGF.EmitBranch(ContBlock);
102 CGF.EmitBlock(ContBlock, true);
104 void Exit(CodeGenFunction &CGF) override {
105 CGF.EmitRuntimeCall(ExitCallee, ExitArgs);
109 // A class to track the execution mode when codegening directives within
110 // a target region. The appropriate mode (generic/spmd) is set on entry
111 // to the target region and used by containing directives such as 'parallel'
112 // to emit optimized code.
113 class ExecutionModeRAII {
115 CGOpenMPRuntimeNVPTX::ExecutionMode SavedMode;
116 CGOpenMPRuntimeNVPTX::ExecutionMode &Mode;
119 ExecutionModeRAII(CGOpenMPRuntimeNVPTX::ExecutionMode &Mode,
120 CGOpenMPRuntimeNVPTX::ExecutionMode NewMode)
125 ~ExecutionModeRAII() { Mode = SavedMode; }
128 /// GPU Configuration: This information can be derived from cuda registers,
129 /// however, providing compile time constants helps generate more efficient
130 /// code. For all practical purposes this is fine because the configuration
131 /// is the same for all known NVPTX architectures.
132 enum MachineConfiguration : unsigned {
134 /// Number of bits required to represent a lane identifier, which is
135 /// computed as log_2(WarpSize).
137 LaneIDMask = WarpSize - 1,
139 /// Global memory alignment for performance.
140 GlobalMemoryAlignment = 256,
143 enum NamedBarrier : unsigned {
144 /// Synchronize on this barrier #ID using a named barrier primitive.
145 /// Only the subset of active threads in a parallel region arrive at the
149 } // anonymous namespace
151 /// Get the GPU warp size.
152 static llvm::Value *getNVPTXWarpSize(CodeGenFunction &CGF) {
153 CGBuilderTy &Bld = CGF.Builder;
154 return Bld.CreateCall(
155 llvm::Intrinsic::getDeclaration(
156 &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_warpsize),
157 llvm::None, "nvptx_warp_size");
160 /// Get the id of the current thread on the GPU.
161 static llvm::Value *getNVPTXThreadID(CodeGenFunction &CGF) {
162 CGBuilderTy &Bld = CGF.Builder;
163 return Bld.CreateCall(
164 llvm::Intrinsic::getDeclaration(
165 &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_tid_x),
166 llvm::None, "nvptx_tid");
169 /// Get the id of the warp in the block.
170 /// We assume that the warp size is 32, which is always the case
171 /// on the NVPTX device, to generate more efficient code.
172 static llvm::Value *getNVPTXWarpID(CodeGenFunction &CGF) {
173 CGBuilderTy &Bld = CGF.Builder;
174 return Bld.CreateAShr(getNVPTXThreadID(CGF), LaneIDBits, "nvptx_warp_id");
177 /// Get the id of the current lane in the Warp.
178 /// We assume that the warp size is 32, which is always the case
179 /// on the NVPTX device, to generate more efficient code.
180 static llvm::Value *getNVPTXLaneID(CodeGenFunction &CGF) {
181 CGBuilderTy &Bld = CGF.Builder;
182 return Bld.CreateAnd(getNVPTXThreadID(CGF), Bld.getInt32(LaneIDMask),
186 /// Get the maximum number of threads in a block of the GPU.
187 static llvm::Value *getNVPTXNumThreads(CodeGenFunction &CGF) {
188 CGBuilderTy &Bld = CGF.Builder;
189 return Bld.CreateCall(
190 llvm::Intrinsic::getDeclaration(
191 &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_ntid_x),
192 llvm::None, "nvptx_num_threads");
195 /// Get barrier to synchronize all threads in a block.
196 static void getNVPTXCTABarrier(CodeGenFunction &CGF) {
197 CGBuilderTy &Bld = CGF.Builder;
198 Bld.CreateCall(llvm::Intrinsic::getDeclaration(
199 &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_barrier0));
202 /// Get barrier #ID to synchronize selected (multiple of warp size) threads in
204 static void getNVPTXBarrier(CodeGenFunction &CGF, int ID,
205 llvm::Value *NumThreads) {
206 CGBuilderTy &Bld = CGF.Builder;
207 llvm::Value *Args[] = {Bld.getInt32(ID), NumThreads};
208 Bld.CreateCall(llvm::Intrinsic::getDeclaration(&CGF.CGM.getModule(),
209 llvm::Intrinsic::nvvm_barrier),
213 /// Synchronize all GPU threads in a block.
214 static void syncCTAThreads(CodeGenFunction &CGF) { getNVPTXCTABarrier(CGF); }
216 /// Synchronize worker threads in a parallel region.
217 static void syncParallelThreads(CodeGenFunction &CGF, llvm::Value *NumThreads) {
218 return getNVPTXBarrier(CGF, NB_Parallel, NumThreads);
221 /// Get the value of the thread_limit clause in the teams directive.
222 /// For the 'generic' execution mode, the runtime encodes thread_limit in
223 /// the launch parameters, always starting thread_limit+warpSize threads per
224 /// CTA. The threads in the last warp are reserved for master execution.
225 /// For the 'spmd' execution mode, all threads in a CTA are part of the team.
226 static llvm::Value *getThreadLimit(CodeGenFunction &CGF,
227 bool IsInSpmdExecutionMode = false) {
228 CGBuilderTy &Bld = CGF.Builder;
229 return IsInSpmdExecutionMode
230 ? getNVPTXNumThreads(CGF)
231 : Bld.CreateSub(getNVPTXNumThreads(CGF), getNVPTXWarpSize(CGF),
235 /// Get the thread id of the OMP master thread.
236 /// The master thread id is the first thread (lane) of the last warp in the
237 /// GPU block. Warp size is assumed to be some power of 2.
238 /// Thread id is 0 indexed.
239 /// E.g: If NumThreads is 33, master id is 32.
240 /// If NumThreads is 64, master id is 32.
241 /// If NumThreads is 1024, master id is 992.
242 static llvm::Value *getMasterThreadID(CodeGenFunction &CGF) {
243 CGBuilderTy &Bld = CGF.Builder;
244 llvm::Value *NumThreads = getNVPTXNumThreads(CGF);
246 // We assume that the warp size is a power of 2.
247 llvm::Value *Mask = Bld.CreateSub(getNVPTXWarpSize(CGF), Bld.getInt32(1));
249 return Bld.CreateAnd(Bld.CreateSub(NumThreads, Bld.getInt32(1)),
250 Bld.CreateNot(Mask), "master_tid");
253 CGOpenMPRuntimeNVPTX::WorkerFunctionState::WorkerFunctionState(
255 : WorkerFn(nullptr), CGFI(nullptr) {
256 createWorkerFunction(CGM);
259 void CGOpenMPRuntimeNVPTX::WorkerFunctionState::createWorkerFunction(
260 CodeGenModule &CGM) {
261 // Create an worker function with no arguments.
262 CGFI = &CGM.getTypes().arrangeNullaryFunction();
264 WorkerFn = llvm::Function::Create(
265 CGM.getTypes().GetFunctionType(*CGFI), llvm::GlobalValue::InternalLinkage,
266 /* placeholder */ "_worker", &CGM.getModule());
267 CGM.SetInternalFunctionAttributes(/*D=*/nullptr, WorkerFn, *CGFI);
270 bool CGOpenMPRuntimeNVPTX::isInSpmdExecutionMode() const {
271 return CurrentExecutionMode == CGOpenMPRuntimeNVPTX::ExecutionMode::Spmd;
274 static CGOpenMPRuntimeNVPTX::ExecutionMode
275 getExecutionModeForDirective(CodeGenModule &CGM,
276 const OMPExecutableDirective &D) {
277 OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
278 switch (DirectiveKind) {
280 case OMPD_target_teams:
281 return CGOpenMPRuntimeNVPTX::ExecutionMode::Generic;
282 case OMPD_target_parallel:
283 return CGOpenMPRuntimeNVPTX::ExecutionMode::Spmd;
285 llvm_unreachable("Unsupported directive on NVPTX device.");
287 llvm_unreachable("Unsupported directive on NVPTX device.");
290 void CGOpenMPRuntimeNVPTX::emitGenericKernel(const OMPExecutableDirective &D,
291 StringRef ParentName,
292 llvm::Function *&OutlinedFn,
293 llvm::Constant *&OutlinedFnID,
295 const RegionCodeGenTy &CodeGen) {
296 ExecutionModeRAII ModeRAII(CurrentExecutionMode,
297 CGOpenMPRuntimeNVPTX::ExecutionMode::Generic);
298 EntryFunctionState EST;
299 WorkerFunctionState WST(CGM);
302 // Emit target region as a standalone region.
303 class NVPTXPrePostActionTy : public PrePostActionTy {
304 CGOpenMPRuntimeNVPTX &RT;
305 CGOpenMPRuntimeNVPTX::EntryFunctionState &EST;
306 CGOpenMPRuntimeNVPTX::WorkerFunctionState &WST;
309 NVPTXPrePostActionTy(CGOpenMPRuntimeNVPTX &RT,
310 CGOpenMPRuntimeNVPTX::EntryFunctionState &EST,
311 CGOpenMPRuntimeNVPTX::WorkerFunctionState &WST)
312 : RT(RT), EST(EST), WST(WST) {}
313 void Enter(CodeGenFunction &CGF) override {
314 RT.emitGenericEntryHeader(CGF, EST, WST);
316 void Exit(CodeGenFunction &CGF) override {
317 RT.emitGenericEntryFooter(CGF, EST);
319 } Action(*this, EST, WST);
320 CodeGen.setAction(Action);
321 emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
322 IsOffloadEntry, CodeGen);
324 // Create the worker function
325 emitWorkerFunction(WST);
327 // Now change the name of the worker function to correspond to this target
328 // region's entry function.
329 WST.WorkerFn->setName(OutlinedFn->getName() + "_worker");
332 // Setup NVPTX threads for master-worker OpenMP scheme.
333 void CGOpenMPRuntimeNVPTX::emitGenericEntryHeader(CodeGenFunction &CGF,
334 EntryFunctionState &EST,
335 WorkerFunctionState &WST) {
336 CGBuilderTy &Bld = CGF.Builder;
338 llvm::BasicBlock *WorkerBB = CGF.createBasicBlock(".worker");
339 llvm::BasicBlock *MasterCheckBB = CGF.createBasicBlock(".mastercheck");
340 llvm::BasicBlock *MasterBB = CGF.createBasicBlock(".master");
341 EST.ExitBB = CGF.createBasicBlock(".exit");
344 Bld.CreateICmpULT(getNVPTXThreadID(CGF), getThreadLimit(CGF));
345 Bld.CreateCondBr(IsWorker, WorkerBB, MasterCheckBB);
347 CGF.EmitBlock(WorkerBB);
348 CGF.EmitCallOrInvoke(WST.WorkerFn, llvm::None);
349 CGF.EmitBranch(EST.ExitBB);
351 CGF.EmitBlock(MasterCheckBB);
353 Bld.CreateICmpEQ(getNVPTXThreadID(CGF), getMasterThreadID(CGF));
354 Bld.CreateCondBr(IsMaster, MasterBB, EST.ExitBB);
356 CGF.EmitBlock(MasterBB);
357 // First action in sequential region:
358 // Initialize the state of the OpenMP runtime library on the GPU.
359 llvm::Value *Args[] = {getThreadLimit(CGF)};
361 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_init), Args);
364 void CGOpenMPRuntimeNVPTX::emitGenericEntryFooter(CodeGenFunction &CGF,
365 EntryFunctionState &EST) {
367 EST.ExitBB = CGF.createBasicBlock(".exit");
369 llvm::BasicBlock *TerminateBB = CGF.createBasicBlock(".termination.notifier");
370 CGF.EmitBranch(TerminateBB);
372 CGF.EmitBlock(TerminateBB);
373 // Signal termination condition.
375 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_deinit), None);
376 // Barrier to terminate worker threads.
378 // Master thread jumps to exit point.
379 CGF.EmitBranch(EST.ExitBB);
381 CGF.EmitBlock(EST.ExitBB);
382 EST.ExitBB = nullptr;
385 void CGOpenMPRuntimeNVPTX::emitSpmdKernel(const OMPExecutableDirective &D,
386 StringRef ParentName,
387 llvm::Function *&OutlinedFn,
388 llvm::Constant *&OutlinedFnID,
390 const RegionCodeGenTy &CodeGen) {
391 ExecutionModeRAII ModeRAII(CurrentExecutionMode,
392 CGOpenMPRuntimeNVPTX::ExecutionMode::Spmd);
393 EntryFunctionState EST;
395 // Emit target region as a standalone region.
396 class NVPTXPrePostActionTy : public PrePostActionTy {
397 CGOpenMPRuntimeNVPTX &RT;
398 CGOpenMPRuntimeNVPTX::EntryFunctionState &EST;
399 const OMPExecutableDirective &D;
402 NVPTXPrePostActionTy(CGOpenMPRuntimeNVPTX &RT,
403 CGOpenMPRuntimeNVPTX::EntryFunctionState &EST,
404 const OMPExecutableDirective &D)
405 : RT(RT), EST(EST), D(D) {}
406 void Enter(CodeGenFunction &CGF) override {
407 RT.emitSpmdEntryHeader(CGF, EST, D);
409 void Exit(CodeGenFunction &CGF) override {
410 RT.emitSpmdEntryFooter(CGF, EST);
412 } Action(*this, EST, D);
413 CodeGen.setAction(Action);
414 emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
415 IsOffloadEntry, CodeGen);
419 void CGOpenMPRuntimeNVPTX::emitSpmdEntryHeader(
420 CodeGenFunction &CGF, EntryFunctionState &EST,
421 const OMPExecutableDirective &D) {
422 auto &Bld = CGF.Builder;
424 // Setup BBs in entry function.
425 llvm::BasicBlock *ExecuteBB = CGF.createBasicBlock(".execute");
426 EST.ExitBB = CGF.createBasicBlock(".exit");
428 // Initialize the OMP state in the runtime; called by all active threads.
429 // TODO: Set RequiresOMPRuntime and RequiresDataSharing parameters
430 // based on code analysis of the target region.
431 llvm::Value *Args[] = {getThreadLimit(CGF, /*IsInSpmdExecutionMode=*/true),
432 /*RequiresOMPRuntime=*/Bld.getInt16(1),
433 /*RequiresDataSharing=*/Bld.getInt16(1)};
435 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_spmd_kernel_init), Args);
436 CGF.EmitBranch(ExecuteBB);
438 CGF.EmitBlock(ExecuteBB);
441 void CGOpenMPRuntimeNVPTX::emitSpmdEntryFooter(CodeGenFunction &CGF,
442 EntryFunctionState &EST) {
444 EST.ExitBB = CGF.createBasicBlock(".exit");
446 llvm::BasicBlock *OMPDeInitBB = CGF.createBasicBlock(".omp.deinit");
447 CGF.EmitBranch(OMPDeInitBB);
449 CGF.EmitBlock(OMPDeInitBB);
450 // DeInitialize the OMP state in the runtime; called by all active threads.
452 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_spmd_kernel_deinit), None);
453 CGF.EmitBranch(EST.ExitBB);
455 CGF.EmitBlock(EST.ExitBB);
456 EST.ExitBB = nullptr;
459 // Create a unique global variable to indicate the execution mode of this target
460 // region. The execution mode is either 'generic', or 'spmd' depending on the
461 // target directive. This variable is picked up by the offload library to setup
462 // the device appropriately before kernel launch. If the execution mode is
463 // 'generic', the runtime reserves one warp for the master, otherwise, all
464 // warps participate in parallel work.
465 static void setPropertyExecutionMode(CodeGenModule &CGM, StringRef Name,
466 CGOpenMPRuntimeNVPTX::ExecutionMode Mode) {
467 (void)new llvm::GlobalVariable(
468 CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
469 llvm::GlobalValue::WeakAnyLinkage,
470 llvm::ConstantInt::get(CGM.Int8Ty, Mode), Name + Twine("_exec_mode"));
473 void CGOpenMPRuntimeNVPTX::emitWorkerFunction(WorkerFunctionState &WST) {
474 auto &Ctx = CGM.getContext();
476 CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
477 CGF.disableDebugInfo();
478 CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, WST.WorkerFn, *WST.CGFI, {});
479 emitWorkerLoop(CGF, WST);
480 CGF.FinishFunction();
483 void CGOpenMPRuntimeNVPTX::emitWorkerLoop(CodeGenFunction &CGF,
484 WorkerFunctionState &WST) {
486 // The workers enter this loop and wait for parallel work from the master.
487 // When the master encounters a parallel region it sets up the work + variable
488 // arguments, and wakes up the workers. The workers first check to see if
489 // they are required for the parallel region, i.e., within the # of requested
490 // parallel threads. The activated workers load the variable arguments and
491 // execute the parallel work.
494 CGBuilderTy &Bld = CGF.Builder;
496 llvm::BasicBlock *AwaitBB = CGF.createBasicBlock(".await.work");
497 llvm::BasicBlock *SelectWorkersBB = CGF.createBasicBlock(".select.workers");
498 llvm::BasicBlock *ExecuteBB = CGF.createBasicBlock(".execute.parallel");
499 llvm::BasicBlock *TerminateBB = CGF.createBasicBlock(".terminate.parallel");
500 llvm::BasicBlock *BarrierBB = CGF.createBasicBlock(".barrier.parallel");
501 llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".exit");
503 CGF.EmitBranch(AwaitBB);
505 // Workers wait for work from master.
506 CGF.EmitBlock(AwaitBB);
507 // Wait for parallel work
511 CGF.CreateDefaultAlignTempAlloca(CGF.Int8PtrTy, /*Name=*/"work_fn");
513 CGF.CreateDefaultAlignTempAlloca(CGF.Int8Ty, /*Name=*/"exec_status");
514 CGF.InitTempAlloca(ExecStatus, Bld.getInt8(/*C=*/0));
515 CGF.InitTempAlloca(WorkFn, llvm::Constant::getNullValue(CGF.Int8PtrTy));
517 llvm::Value *Args[] = {WorkFn.getPointer()};
518 llvm::Value *Ret = CGF.EmitRuntimeCall(
519 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_parallel), Args);
520 Bld.CreateStore(Bld.CreateZExt(Ret, CGF.Int8Ty), ExecStatus);
522 // On termination condition (workid == 0), exit loop.
523 llvm::Value *ShouldTerminate =
524 Bld.CreateIsNull(Bld.CreateLoad(WorkFn), "should_terminate");
525 Bld.CreateCondBr(ShouldTerminate, ExitBB, SelectWorkersBB);
527 // Activate requested workers.
528 CGF.EmitBlock(SelectWorkersBB);
529 llvm::Value *IsActive =
530 Bld.CreateIsNotNull(Bld.CreateLoad(ExecStatus), "is_active");
531 Bld.CreateCondBr(IsActive, ExecuteBB, BarrierBB);
533 // Signal start of parallel region.
534 CGF.EmitBlock(ExecuteBB);
536 // Process work items: outlined parallel functions.
537 for (auto *W : Work) {
538 // Try to match this outlined function.
539 auto *ID = Bld.CreatePointerBitCastOrAddrSpaceCast(W, CGM.Int8PtrTy);
541 llvm::Value *WorkFnMatch =
542 Bld.CreateICmpEQ(Bld.CreateLoad(WorkFn), ID, "work_match");
544 llvm::BasicBlock *ExecuteFNBB = CGF.createBasicBlock(".execute.fn");
545 llvm::BasicBlock *CheckNextBB = CGF.createBasicBlock(".check.next");
546 Bld.CreateCondBr(WorkFnMatch, ExecuteFNBB, CheckNextBB);
548 // Execute this outlined function.
549 CGF.EmitBlock(ExecuteFNBB);
551 // Insert call to work function.
552 // FIXME: Pass arguments to outlined function from master thread.
553 auto *Fn = cast<llvm::Function>(W);
555 CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty, /*Name=*/".zero.addr");
556 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C=*/0));
557 llvm::Value *FnArgs[] = {ZeroAddr.getPointer(), ZeroAddr.getPointer()};
558 CGF.EmitCallOrInvoke(Fn, FnArgs);
560 // Go to end of parallel region.
561 CGF.EmitBranch(TerminateBB);
563 CGF.EmitBlock(CheckNextBB);
566 // Signal end of parallel region.
567 CGF.EmitBlock(TerminateBB);
569 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_end_parallel),
571 CGF.EmitBranch(BarrierBB);
573 // All active and inactive workers wait at a barrier after parallel region.
574 CGF.EmitBlock(BarrierBB);
575 // Barrier after parallel region.
577 CGF.EmitBranch(AwaitBB);
579 // Exit target region.
580 CGF.EmitBlock(ExitBB);
583 /// \brief Returns specified OpenMP runtime function for the current OpenMP
584 /// implementation. Specialized for the NVPTX device.
585 /// \param Function OpenMP runtime function.
586 /// \return Specified function.
588 CGOpenMPRuntimeNVPTX::createNVPTXRuntimeFunction(unsigned Function) {
589 llvm::Constant *RTLFn = nullptr;
590 switch (static_cast<OpenMPRTLFunctionNVPTX>(Function)) {
591 case OMPRTL_NVPTX__kmpc_kernel_init: {
592 // Build void __kmpc_kernel_init(kmp_int32 thread_limit);
593 llvm::Type *TypeParams[] = {CGM.Int32Ty};
594 llvm::FunctionType *FnTy =
595 llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
596 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_init");
599 case OMPRTL_NVPTX__kmpc_kernel_deinit: {
600 // Build void __kmpc_kernel_deinit();
601 llvm::FunctionType *FnTy =
602 llvm::FunctionType::get(CGM.VoidTy, llvm::None, /*isVarArg*/ false);
603 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_deinit");
606 case OMPRTL_NVPTX__kmpc_spmd_kernel_init: {
607 // Build void __kmpc_spmd_kernel_init(kmp_int32 thread_limit,
608 // short RequiresOMPRuntime, short RequiresDataSharing);
609 llvm::Type *TypeParams[] = {CGM.Int32Ty, CGM.Int16Ty, CGM.Int16Ty};
610 llvm::FunctionType *FnTy =
611 llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
612 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_spmd_kernel_init");
615 case OMPRTL_NVPTX__kmpc_spmd_kernel_deinit: {
616 // Build void __kmpc_spmd_kernel_deinit();
617 llvm::FunctionType *FnTy =
618 llvm::FunctionType::get(CGM.VoidTy, llvm::None, /*isVarArg*/ false);
619 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_spmd_kernel_deinit");
622 case OMPRTL_NVPTX__kmpc_kernel_prepare_parallel: {
623 /// Build void __kmpc_kernel_prepare_parallel(
624 /// void *outlined_function);
625 llvm::Type *TypeParams[] = {CGM.Int8PtrTy};
626 llvm::FunctionType *FnTy =
627 llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
628 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_prepare_parallel");
631 case OMPRTL_NVPTX__kmpc_kernel_parallel: {
632 /// Build bool __kmpc_kernel_parallel(void **outlined_function);
633 llvm::Type *TypeParams[] = {CGM.Int8PtrPtrTy};
634 llvm::Type *RetTy = CGM.getTypes().ConvertType(CGM.getContext().BoolTy);
635 llvm::FunctionType *FnTy =
636 llvm::FunctionType::get(RetTy, TypeParams, /*isVarArg*/ false);
637 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_parallel");
640 case OMPRTL_NVPTX__kmpc_kernel_end_parallel: {
641 /// Build void __kmpc_kernel_end_parallel();
642 llvm::FunctionType *FnTy =
643 llvm::FunctionType::get(CGM.VoidTy, llvm::None, /*isVarArg*/ false);
644 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_end_parallel");
647 case OMPRTL_NVPTX__kmpc_serialized_parallel: {
648 // Build void __kmpc_serialized_parallel(ident_t *loc, kmp_int32
650 llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty};
651 llvm::FunctionType *FnTy =
652 llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
653 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_serialized_parallel");
656 case OMPRTL_NVPTX__kmpc_end_serialized_parallel: {
657 // Build void __kmpc_end_serialized_parallel(ident_t *loc, kmp_int32
659 llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty};
660 llvm::FunctionType *FnTy =
661 llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
662 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_end_serialized_parallel");
665 case OMPRTL_NVPTX__kmpc_shuffle_int32: {
666 // Build int32_t __kmpc_shuffle_int32(int32_t element,
667 // int16_t lane_offset, int16_t warp_size);
668 llvm::Type *TypeParams[] = {CGM.Int32Ty, CGM.Int16Ty, CGM.Int16Ty};
669 llvm::FunctionType *FnTy =
670 llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
671 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_shuffle_int32");
674 case OMPRTL_NVPTX__kmpc_shuffle_int64: {
675 // Build int64_t __kmpc_shuffle_int64(int64_t element,
676 // int16_t lane_offset, int16_t warp_size);
677 llvm::Type *TypeParams[] = {CGM.Int64Ty, CGM.Int16Ty, CGM.Int16Ty};
678 llvm::FunctionType *FnTy =
679 llvm::FunctionType::get(CGM.Int64Ty, TypeParams, /*isVarArg*/ false);
680 RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_shuffle_int64");
683 case OMPRTL_NVPTX__kmpc_parallel_reduce_nowait: {
684 // Build int32_t kmpc_nvptx_parallel_reduce_nowait(kmp_int32 global_tid,
685 // kmp_int32 num_vars, size_t reduce_size, void* reduce_data,
686 // void (*kmp_ShuffleReductFctPtr)(void *rhsData, int16_t lane_id, int16_t
687 // lane_offset, int16_t Algorithm Version),
688 // void (*kmp_InterWarpCopyFctPtr)(void* src, int warp_num));
689 llvm::Type *ShuffleReduceTypeParams[] = {CGM.VoidPtrTy, CGM.Int16Ty,
690 CGM.Int16Ty, CGM.Int16Ty};
691 auto *ShuffleReduceFnTy =
692 llvm::FunctionType::get(CGM.VoidTy, ShuffleReduceTypeParams,
694 llvm::Type *InterWarpCopyTypeParams[] = {CGM.VoidPtrTy, CGM.Int32Ty};
695 auto *InterWarpCopyFnTy =
696 llvm::FunctionType::get(CGM.VoidTy, InterWarpCopyTypeParams,
698 llvm::Type *TypeParams[] = {CGM.Int32Ty,
702 ShuffleReduceFnTy->getPointerTo(),
703 InterWarpCopyFnTy->getPointerTo()};
704 llvm::FunctionType *FnTy =
705 llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg=*/false);
706 RTLFn = CGM.CreateRuntimeFunction(
707 FnTy, /*Name=*/"__kmpc_nvptx_parallel_reduce_nowait");
710 case OMPRTL_NVPTX__kmpc_teams_reduce_nowait: {
711 // Build int32_t __kmpc_nvptx_teams_reduce_nowait(int32_t global_tid,
712 // int32_t num_vars, size_t reduce_size, void *reduce_data,
713 // void (*kmp_ShuffleReductFctPtr)(void *rhsData, int16_t lane_id, int16_t
714 // lane_offset, int16_t shortCircuit),
715 // void (*kmp_InterWarpCopyFctPtr)(void* src, int32_t warp_num),
716 // void (*kmp_CopyToScratchpadFctPtr)(void *reduce_data, void * scratchpad,
717 // int32_t index, int32_t width),
718 // void (*kmp_LoadReduceFctPtr)(void *reduce_data, void * scratchpad,
719 // int32_t index, int32_t width, int32_t reduce))
720 llvm::Type *ShuffleReduceTypeParams[] = {CGM.VoidPtrTy, CGM.Int16Ty,
721 CGM.Int16Ty, CGM.Int16Ty};
722 auto *ShuffleReduceFnTy =
723 llvm::FunctionType::get(CGM.VoidTy, ShuffleReduceTypeParams,
725 llvm::Type *InterWarpCopyTypeParams[] = {CGM.VoidPtrTy, CGM.Int32Ty};
726 auto *InterWarpCopyFnTy =
727 llvm::FunctionType::get(CGM.VoidTy, InterWarpCopyTypeParams,
729 llvm::Type *CopyToScratchpadTypeParams[] = {CGM.VoidPtrTy, CGM.VoidPtrTy,
730 CGM.Int32Ty, CGM.Int32Ty};
731 auto *CopyToScratchpadFnTy =
732 llvm::FunctionType::get(CGM.VoidTy, CopyToScratchpadTypeParams,
734 llvm::Type *LoadReduceTypeParams[] = {
735 CGM.VoidPtrTy, CGM.VoidPtrTy, CGM.Int32Ty, CGM.Int32Ty, CGM.Int32Ty};
736 auto *LoadReduceFnTy =
737 llvm::FunctionType::get(CGM.VoidTy, LoadReduceTypeParams,
739 llvm::Type *TypeParams[] = {CGM.Int32Ty,
743 ShuffleReduceFnTy->getPointerTo(),
744 InterWarpCopyFnTy->getPointerTo(),
745 CopyToScratchpadFnTy->getPointerTo(),
746 LoadReduceFnTy->getPointerTo()};
747 llvm::FunctionType *FnTy =
748 llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg=*/false);
749 RTLFn = CGM.CreateRuntimeFunction(
750 FnTy, /*Name=*/"__kmpc_nvptx_teams_reduce_nowait");
753 case OMPRTL_NVPTX__kmpc_end_reduce_nowait: {
754 // Build __kmpc_end_reduce_nowait(kmp_int32 global_tid);
755 llvm::Type *TypeParams[] = {CGM.Int32Ty};
756 llvm::FunctionType *FnTy =
757 llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
758 RTLFn = CGM.CreateRuntimeFunction(
759 FnTy, /*Name=*/"__kmpc_nvptx_end_reduce_nowait");
766 void CGOpenMPRuntimeNVPTX::createOffloadEntry(llvm::Constant *ID,
767 llvm::Constant *Addr,
768 uint64_t Size, int32_t) {
769 auto *F = dyn_cast<llvm::Function>(Addr);
770 // TODO: Add support for global variables on the device after declare target
774 llvm::Module *M = F->getParent();
775 llvm::LLVMContext &Ctx = M->getContext();
777 // Get "nvvm.annotations" metadata node
778 llvm::NamedMDNode *MD = M->getOrInsertNamedMetadata("nvvm.annotations");
780 llvm::Metadata *MDVals[] = {
781 llvm::ConstantAsMetadata::get(F), llvm::MDString::get(Ctx, "kernel"),
782 llvm::ConstantAsMetadata::get(
783 llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), 1))};
784 // Append metadata to nvvm.annotations
785 MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
788 void CGOpenMPRuntimeNVPTX::emitTargetOutlinedFunction(
789 const OMPExecutableDirective &D, StringRef ParentName,
790 llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
791 bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
792 if (!IsOffloadEntry) // Nothing to do.
795 assert(!ParentName.empty() && "Invalid target region parent name!");
797 CGOpenMPRuntimeNVPTX::ExecutionMode Mode =
798 getExecutionModeForDirective(CGM, D);
800 case CGOpenMPRuntimeNVPTX::ExecutionMode::Generic:
801 emitGenericKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
804 case CGOpenMPRuntimeNVPTX::ExecutionMode::Spmd:
805 emitSpmdKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
808 case CGOpenMPRuntimeNVPTX::ExecutionMode::Unknown:
810 "Unknown programming model for OpenMP directive on NVPTX target.");
813 setPropertyExecutionMode(CGM, OutlinedFn->getName(), Mode);
816 CGOpenMPRuntimeNVPTX::CGOpenMPRuntimeNVPTX(CodeGenModule &CGM)
817 : CGOpenMPRuntime(CGM), CurrentExecutionMode(ExecutionMode::Unknown) {
818 if (!CGM.getLangOpts().OpenMPIsDevice)
819 llvm_unreachable("OpenMP NVPTX can only handle device code.");
822 void CGOpenMPRuntimeNVPTX::emitProcBindClause(CodeGenFunction &CGF,
823 OpenMPProcBindClauseKind ProcBind,
824 SourceLocation Loc) {
825 // Do nothing in case of Spmd mode and L0 parallel.
826 // TODO: If in Spmd mode and L1 parallel emit the clause.
827 if (isInSpmdExecutionMode())
830 CGOpenMPRuntime::emitProcBindClause(CGF, ProcBind, Loc);
833 void CGOpenMPRuntimeNVPTX::emitNumThreadsClause(CodeGenFunction &CGF,
834 llvm::Value *NumThreads,
835 SourceLocation Loc) {
836 // Do nothing in case of Spmd mode and L0 parallel.
837 // TODO: If in Spmd mode and L1 parallel emit the clause.
838 if (isInSpmdExecutionMode())
841 CGOpenMPRuntime::emitNumThreadsClause(CGF, NumThreads, Loc);
844 void CGOpenMPRuntimeNVPTX::emitNumTeamsClause(CodeGenFunction &CGF,
845 const Expr *NumTeams,
846 const Expr *ThreadLimit,
847 SourceLocation Loc) {}
849 llvm::Value *CGOpenMPRuntimeNVPTX::emitParallelOutlinedFunction(
850 const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
851 OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
852 return CGOpenMPRuntime::emitParallelOutlinedFunction(D, ThreadIDVar,
853 InnermostKind, CodeGen);
856 llvm::Value *CGOpenMPRuntimeNVPTX::emitTeamsOutlinedFunction(
857 const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
858 OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
860 llvm::Value *OutlinedFunVal = CGOpenMPRuntime::emitTeamsOutlinedFunction(
861 D, ThreadIDVar, InnermostKind, CodeGen);
862 llvm::Function *OutlinedFun = cast<llvm::Function>(OutlinedFunVal);
863 OutlinedFun->removeFnAttr(llvm::Attribute::NoInline);
864 OutlinedFun->addFnAttr(llvm::Attribute::AlwaysInline);
869 void CGOpenMPRuntimeNVPTX::emitTeamsCall(CodeGenFunction &CGF,
870 const OMPExecutableDirective &D,
872 llvm::Value *OutlinedFn,
873 ArrayRef<llvm::Value *> CapturedVars) {
874 if (!CGF.HaveInsertPoint())
878 CGF.CreateTempAlloca(CGF.Int32Ty, CharUnits::fromQuantity(4),
879 /*Name*/ ".zero.addr");
880 CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
881 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
882 OutlinedFnArgs.push_back(ZeroAddr.getPointer());
883 OutlinedFnArgs.push_back(ZeroAddr.getPointer());
884 OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
885 CGF.EmitCallOrInvoke(OutlinedFn, OutlinedFnArgs);
888 void CGOpenMPRuntimeNVPTX::emitParallelCall(
889 CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
890 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
891 if (!CGF.HaveInsertPoint())
894 if (isInSpmdExecutionMode())
895 emitSpmdParallelCall(CGF, Loc, OutlinedFn, CapturedVars, IfCond);
897 emitGenericParallelCall(CGF, Loc, OutlinedFn, CapturedVars, IfCond);
900 void CGOpenMPRuntimeNVPTX::emitGenericParallelCall(
901 CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
902 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
903 llvm::Function *Fn = cast<llvm::Function>(OutlinedFn);
905 auto &&L0ParallelGen = [this, Fn](CodeGenFunction &CGF, PrePostActionTy &) {
906 CGBuilderTy &Bld = CGF.Builder;
908 // Prepare for parallel region. Indicate the outlined function.
909 llvm::Value *Args[] = {Bld.CreateBitOrPointerCast(Fn, CGM.Int8PtrTy)};
911 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_prepare_parallel),
914 // Activate workers. This barrier is used by the master to signal
915 // work for the workers.
918 // OpenMP [2.5, Parallel Construct, p.49]
919 // There is an implied barrier at the end of a parallel region. After the
920 // end of a parallel region, only the master thread of the team resumes
921 // execution of the enclosing task region.
923 // The master waits at this barrier until all workers are done.
926 // Remember for post-processing in worker loop.
930 auto *RTLoc = emitUpdateLocation(CGF, Loc);
931 auto *ThreadID = getThreadID(CGF, Loc);
932 llvm::Value *Args[] = {RTLoc, ThreadID};
934 auto &&SeqGen = [this, Fn, &CapturedVars, &Args](CodeGenFunction &CGF,
936 auto &&CodeGen = [this, Fn, &CapturedVars](CodeGenFunction &CGF,
937 PrePostActionTy &Action) {
940 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
941 OutlinedFnArgs.push_back(
942 llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
943 OutlinedFnArgs.push_back(
944 llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
945 OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
946 CGF.EmitCallOrInvoke(Fn, OutlinedFnArgs);
949 RegionCodeGenTy RCG(CodeGen);
950 NVPTXActionTy Action(
951 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_serialized_parallel),
953 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_end_serialized_parallel),
955 RCG.setAction(Action);
960 emitOMPIfClause(CGF, IfCond, L0ParallelGen, SeqGen);
962 CodeGenFunction::RunCleanupsScope Scope(CGF);
963 RegionCodeGenTy ThenRCG(L0ParallelGen);
968 void CGOpenMPRuntimeNVPTX::emitSpmdParallelCall(
969 CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
970 ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
971 // Just call the outlined function to execute the parallel region.
972 // OutlinedFn(>id, &zero, CapturedStruct);
974 // TODO: Do something with IfCond when support for the 'if' clause
975 // is added on Spmd target directives.
976 llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
977 OutlinedFnArgs.push_back(
978 llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
979 OutlinedFnArgs.push_back(
980 llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
981 OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
982 CGF.EmitCallOrInvoke(OutlinedFn, OutlinedFnArgs);
985 /// This function creates calls to one of two shuffle functions to copy
986 /// variables between lanes in a warp.
987 static llvm::Value *createRuntimeShuffleFunction(CodeGenFunction &CGF,
990 llvm::Value *Offset) {
992 auto &C = CGM.getContext();
993 auto &Bld = CGF.Builder;
994 CGOpenMPRuntimeNVPTX &RT =
995 *(static_cast<CGOpenMPRuntimeNVPTX *>(&CGM.getOpenMPRuntime()));
997 unsigned Size = CGM.getContext().getTypeSizeInChars(ElemTy).getQuantity();
998 assert(Size <= 8 && "Unsupported bitwidth in shuffle instruction.");
1000 OpenMPRTLFunctionNVPTX ShuffleFn = Size <= 4
1001 ? OMPRTL_NVPTX__kmpc_shuffle_int32
1002 : OMPRTL_NVPTX__kmpc_shuffle_int64;
1004 // Cast all types to 32- or 64-bit values before calling shuffle routines.
1005 auto CastTy = Size <= 4 ? CGM.Int32Ty : CGM.Int64Ty;
1006 auto *ElemCast = Bld.CreateSExtOrBitCast(Elem, CastTy);
1007 auto *WarpSize = CGF.EmitScalarConversion(
1008 getNVPTXWarpSize(CGF), C.getIntTypeForBitwidth(32, /* Signed */ true),
1009 C.getIntTypeForBitwidth(16, /* Signed */ true), SourceLocation());
1012 CGF.EmitRuntimeCall(RT.createNVPTXRuntimeFunction(ShuffleFn),
1013 {ElemCast, Offset, WarpSize});
1015 return Bld.CreateTruncOrBitCast(ShuffledVal, CGF.ConvertTypeForMem(ElemTy));
1019 enum CopyAction : unsigned {
1020 // RemoteLaneToThread: Copy over a Reduce list from a remote lane in
1021 // the warp using shuffle instructions.
1023 // ThreadCopy: Make a copy of a Reduce list on the thread's stack.
1025 // ThreadToScratchpad: Copy a team-reduced array to the scratchpad.
1027 // ScratchpadToThread: Copy from a scratchpad array in global memory
1028 // containing team-reduced data to a thread's stack.
1033 struct CopyOptionsTy {
1034 llvm::Value *RemoteLaneOffset;
1035 llvm::Value *ScratchpadIndex;
1036 llvm::Value *ScratchpadWidth;
1039 /// Emit instructions to copy a Reduce list, which contains partially
1040 /// aggregated values, in the specified direction.
1041 static void emitReductionListCopy(
1042 CopyAction Action, CodeGenFunction &CGF, QualType ReductionArrayTy,
1043 ArrayRef<const Expr *> Privates, Address SrcBase, Address DestBase,
1044 CopyOptionsTy CopyOptions = {nullptr, nullptr, nullptr}) {
1046 auto &CGM = CGF.CGM;
1047 auto &C = CGM.getContext();
1048 auto &Bld = CGF.Builder;
1050 auto *RemoteLaneOffset = CopyOptions.RemoteLaneOffset;
1051 auto *ScratchpadIndex = CopyOptions.ScratchpadIndex;
1052 auto *ScratchpadWidth = CopyOptions.ScratchpadWidth;
1054 // Iterates, element-by-element, through the source Reduce list and
1057 unsigned Size = Privates.size();
1058 for (auto &Private : Privates) {
1059 Address SrcElementAddr = Address::invalid();
1060 Address DestElementAddr = Address::invalid();
1061 Address DestElementPtrAddr = Address::invalid();
1062 // Should we shuffle in an element from a remote lane?
1063 bool ShuffleInElement = false;
1064 // Set to true to update the pointer in the dest Reduce list to a
1065 // newly created element.
1066 bool UpdateDestListPtr = false;
1067 // Increment the src or dest pointer to the scratchpad, for each
1069 bool IncrScratchpadSrc = false;
1070 bool IncrScratchpadDest = false;
1073 case RemoteLaneToThread: {
1074 // Step 1.1: Get the address for the src element in the Reduce list.
1075 Address SrcElementPtrAddr =
1076 Bld.CreateConstArrayGEP(SrcBase, Idx, CGF.getPointerSize());
1077 llvm::Value *SrcElementPtrPtr = CGF.EmitLoadOfScalar(
1078 SrcElementPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
1080 Address(SrcElementPtrPtr, C.getTypeAlignInChars(Private->getType()));
1082 // Step 1.2: Create a temporary to store the element in the destination
1084 DestElementPtrAddr =
1085 Bld.CreateConstArrayGEP(DestBase, Idx, CGF.getPointerSize());
1087 CGF.CreateMemTemp(Private->getType(), ".omp.reduction.element");
1088 ShuffleInElement = true;
1089 UpdateDestListPtr = true;
1093 // Step 1.1: Get the address for the src element in the Reduce list.
1094 Address SrcElementPtrAddr =
1095 Bld.CreateConstArrayGEP(SrcBase, Idx, CGF.getPointerSize());
1096 llvm::Value *SrcElementPtrPtr = CGF.EmitLoadOfScalar(
1097 SrcElementPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
1099 Address(SrcElementPtrPtr, C.getTypeAlignInChars(Private->getType()));
1101 // Step 1.2: Get the address for dest element. The destination
1102 // element has already been created on the thread's stack.
1103 DestElementPtrAddr =
1104 Bld.CreateConstArrayGEP(DestBase, Idx, CGF.getPointerSize());
1105 llvm::Value *DestElementPtr =
1106 CGF.EmitLoadOfScalar(DestElementPtrAddr, /*Volatile=*/false,
1107 C.VoidPtrTy, SourceLocation());
1108 Address DestElemAddr =
1109 Address(DestElementPtr, C.getTypeAlignInChars(Private->getType()));
1110 DestElementAddr = Bld.CreateElementBitCast(
1111 DestElemAddr, CGF.ConvertTypeForMem(Private->getType()));
1114 case ThreadToScratchpad: {
1115 // Step 1.1: Get the address for the src element in the Reduce list.
1116 Address SrcElementPtrAddr =
1117 Bld.CreateConstArrayGEP(SrcBase, Idx, CGF.getPointerSize());
1118 llvm::Value *SrcElementPtrPtr = CGF.EmitLoadOfScalar(
1119 SrcElementPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
1121 Address(SrcElementPtrPtr, C.getTypeAlignInChars(Private->getType()));
1123 // Step 1.2: Get the address for dest element:
1124 // address = base + index * ElementSizeInChars.
1125 unsigned ElementSizeInChars =
1126 C.getTypeSizeInChars(Private->getType()).getQuantity();
1127 auto *CurrentOffset =
1128 Bld.CreateMul(llvm::ConstantInt::get(CGM.SizeTy, ElementSizeInChars),
1130 auto *ScratchPadElemAbsolutePtrVal =
1131 Bld.CreateAdd(DestBase.getPointer(), CurrentOffset);
1132 ScratchPadElemAbsolutePtrVal =
1133 Bld.CreateIntToPtr(ScratchPadElemAbsolutePtrVal, CGF.VoidPtrTy);
1134 Address ScratchpadPtr =
1135 Address(ScratchPadElemAbsolutePtrVal,
1136 C.getTypeAlignInChars(Private->getType()));
1137 DestElementAddr = Bld.CreateElementBitCast(
1138 ScratchpadPtr, CGF.ConvertTypeForMem(Private->getType()));
1139 IncrScratchpadDest = true;
1142 case ScratchpadToThread: {
1143 // Step 1.1: Get the address for the src element in the scratchpad.
1144 // address = base + index * ElementSizeInChars.
1145 unsigned ElementSizeInChars =
1146 C.getTypeSizeInChars(Private->getType()).getQuantity();
1147 auto *CurrentOffset =
1148 Bld.CreateMul(llvm::ConstantInt::get(CGM.SizeTy, ElementSizeInChars),
1150 auto *ScratchPadElemAbsolutePtrVal =
1151 Bld.CreateAdd(SrcBase.getPointer(), CurrentOffset);
1152 ScratchPadElemAbsolutePtrVal =
1153 Bld.CreateIntToPtr(ScratchPadElemAbsolutePtrVal, CGF.VoidPtrTy);
1154 SrcElementAddr = Address(ScratchPadElemAbsolutePtrVal,
1155 C.getTypeAlignInChars(Private->getType()));
1156 IncrScratchpadSrc = true;
1158 // Step 1.2: Create a temporary to store the element in the destination
1160 DestElementPtrAddr =
1161 Bld.CreateConstArrayGEP(DestBase, Idx, CGF.getPointerSize());
1163 CGF.CreateMemTemp(Private->getType(), ".omp.reduction.element");
1164 UpdateDestListPtr = true;
1169 // Regardless of src and dest of copy, we emit the load of src
1170 // element as this is required in all directions
1171 SrcElementAddr = Bld.CreateElementBitCast(
1172 SrcElementAddr, CGF.ConvertTypeForMem(Private->getType()));
1174 CGF.EmitLoadOfScalar(SrcElementAddr, /*Volatile=*/false,
1175 Private->getType(), SourceLocation());
1177 // Now that all active lanes have read the element in the
1178 // Reduce list, shuffle over the value from the remote lane.
1179 if (ShuffleInElement) {
1180 Elem = createRuntimeShuffleFunction(CGF, Private->getType(), Elem,
1184 // Store the source element value to the dest element address.
1185 CGF.EmitStoreOfScalar(Elem, DestElementAddr, /*Volatile=*/false,
1186 Private->getType());
1188 // Step 3.1: Modify reference in dest Reduce list as needed.
1189 // Modifying the reference in Reduce list to point to the newly
1190 // created element. The element is live in the current function
1191 // scope and that of functions it invokes (i.e., reduce_function).
1192 // RemoteReduceData[i] = (void*)&RemoteElem
1193 if (UpdateDestListPtr) {
1194 CGF.EmitStoreOfScalar(Bld.CreatePointerBitCastOrAddrSpaceCast(
1195 DestElementAddr.getPointer(), CGF.VoidPtrTy),
1196 DestElementPtrAddr, /*Volatile=*/false,
1200 // Step 4.1: Increment SrcBase/DestBase so that it points to the starting
1201 // address of the next element in scratchpad memory, unless we're currently
1202 // processing the last one. Memory alignment is also taken care of here.
1203 if ((IncrScratchpadDest || IncrScratchpadSrc) && (Idx + 1 < Size)) {
1204 llvm::Value *ScratchpadBasePtr =
1205 IncrScratchpadDest ? DestBase.getPointer() : SrcBase.getPointer();
1206 unsigned ElementSizeInChars =
1207 C.getTypeSizeInChars(Private->getType()).getQuantity();
1208 ScratchpadBasePtr = Bld.CreateAdd(
1210 Bld.CreateMul(ScratchpadWidth, llvm::ConstantInt::get(
1211 CGM.SizeTy, ElementSizeInChars)));
1213 // Take care of global memory alignment for performance
1214 ScratchpadBasePtr = Bld.CreateSub(ScratchpadBasePtr,
1215 llvm::ConstantInt::get(CGM.SizeTy, 1));
1216 ScratchpadBasePtr = Bld.CreateSDiv(
1218 llvm::ConstantInt::get(CGM.SizeTy, GlobalMemoryAlignment));
1219 ScratchpadBasePtr = Bld.CreateAdd(ScratchpadBasePtr,
1220 llvm::ConstantInt::get(CGM.SizeTy, 1));
1221 ScratchpadBasePtr = Bld.CreateMul(
1223 llvm::ConstantInt::get(CGM.SizeTy, GlobalMemoryAlignment));
1225 if (IncrScratchpadDest)
1226 DestBase = Address(ScratchpadBasePtr, CGF.getPointerAlign());
1227 else /* IncrScratchpadSrc = true */
1228 SrcBase = Address(ScratchpadBasePtr, CGF.getPointerAlign());
1235 /// This function emits a helper that loads data from the scratchpad array
1236 /// and (optionally) reduces it with the input operand.
1238 /// load_and_reduce(local, scratchpad, index, width, should_reduce)
1239 /// reduce_data remote;
1240 /// for elem in remote:
1241 /// remote.elem = Scratchpad[elem_id][index]
1242 /// if (should_reduce)
1243 /// local = local @ remote
1246 static llvm::Value *
1247 emitReduceScratchpadFunction(CodeGenModule &CGM,
1248 ArrayRef<const Expr *> Privates,
1249 QualType ReductionArrayTy, llvm::Value *ReduceFn) {
1250 auto &C = CGM.getContext();
1251 auto Int32Ty = C.getIntTypeForBitwidth(32, /* Signed */ true);
1253 // Destination of the copy.
1254 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, SourceLocation(),
1255 /*Id=*/nullptr, C.VoidPtrTy);
1256 // Base address of the scratchpad array, with each element storing a
1257 // Reduce list per team.
1258 ImplicitParamDecl ScratchPadArg(C, /*DC=*/nullptr, SourceLocation(),
1259 /*Id=*/nullptr, C.VoidPtrTy);
1260 // A source index into the scratchpad array.
1261 ImplicitParamDecl IndexArg(C, /*DC=*/nullptr, SourceLocation(),
1262 /*Id=*/nullptr, Int32Ty);
1263 // Row width of an element in the scratchpad array, typically
1264 // the number of teams.
1265 ImplicitParamDecl WidthArg(C, /*DC=*/nullptr, SourceLocation(),
1266 /*Id=*/nullptr, Int32Ty);
1267 // If should_reduce == 1, then it's load AND reduce,
1268 // If should_reduce == 0 (or otherwise), then it only loads (+ copy).
1269 // The latter case is used for initialization.
1270 ImplicitParamDecl ShouldReduceArg(C, /*DC=*/nullptr, SourceLocation(),
1271 /*Id=*/nullptr, Int32Ty);
1273 FunctionArgList Args;
1274 Args.push_back(&ReduceListArg);
1275 Args.push_back(&ScratchPadArg);
1276 Args.push_back(&IndexArg);
1277 Args.push_back(&WidthArg);
1278 Args.push_back(&ShouldReduceArg);
1280 auto &CGFI = CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
1281 auto *Fn = llvm::Function::Create(
1282 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
1283 "_omp_reduction_load_and_reduce", &CGM.getModule());
1284 CGM.SetInternalFunctionAttributes(/*DC=*/nullptr, Fn, CGFI);
1285 CodeGenFunction CGF(CGM);
1286 // We don't need debug information in this function as nothing here refers to
1288 CGF.disableDebugInfo();
1289 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args);
1291 auto &Bld = CGF.Builder;
1293 // Get local Reduce list pointer.
1294 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
1295 Address ReduceListAddr(
1296 Bld.CreatePointerBitCastOrAddrSpaceCast(
1297 CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
1298 C.VoidPtrTy, SourceLocation()),
1299 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
1300 CGF.getPointerAlign());
1302 Address AddrScratchPadArg = CGF.GetAddrOfLocalVar(&ScratchPadArg);
1303 llvm::Value *ScratchPadBase = CGF.EmitLoadOfScalar(
1304 AddrScratchPadArg, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
1306 Address AddrIndexArg = CGF.GetAddrOfLocalVar(&IndexArg);
1307 llvm::Value *IndexVal =
1308 Bld.CreateIntCast(CGF.EmitLoadOfScalar(AddrIndexArg, /*Volatile=*/false,
1309 Int32Ty, SourceLocation()),
1310 CGM.SizeTy, /*isSigned=*/true);
1312 Address AddrWidthArg = CGF.GetAddrOfLocalVar(&WidthArg);
1313 llvm::Value *WidthVal =
1314 Bld.CreateIntCast(CGF.EmitLoadOfScalar(AddrWidthArg, /*Volatile=*/false,
1315 Int32Ty, SourceLocation()),
1316 CGM.SizeTy, /*isSigned=*/true);
1318 Address AddrShouldReduceArg = CGF.GetAddrOfLocalVar(&ShouldReduceArg);
1319 llvm::Value *ShouldReduceVal = CGF.EmitLoadOfScalar(
1320 AddrShouldReduceArg, /*Volatile=*/false, Int32Ty, SourceLocation());
1322 // The absolute ptr address to the base addr of the next element to copy.
1323 llvm::Value *CumulativeElemBasePtr =
1324 Bld.CreatePtrToInt(ScratchPadBase, CGM.SizeTy);
1325 Address SrcDataAddr(CumulativeElemBasePtr, CGF.getPointerAlign());
1327 // Create a Remote Reduce list to store the elements read from the
1328 // scratchpad array.
1329 Address RemoteReduceList =
1330 CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.remote_red_list");
1332 // Assemble remote Reduce list from scratchpad array.
1333 emitReductionListCopy(ScratchpadToThread, CGF, ReductionArrayTy, Privates,
1334 SrcDataAddr, RemoteReduceList,
1335 {/*RemoteLaneOffset=*/nullptr,
1336 /*ScratchpadIndex=*/IndexVal,
1337 /*ScratchpadWidth=*/WidthVal});
1339 llvm::BasicBlock *ThenBB = CGF.createBasicBlock("then");
1340 llvm::BasicBlock *ElseBB = CGF.createBasicBlock("else");
1341 llvm::BasicBlock *MergeBB = CGF.createBasicBlock("ifcont");
1343 auto CondReduce = Bld.CreateICmpEQ(ShouldReduceVal, Bld.getInt32(1));
1344 Bld.CreateCondBr(CondReduce, ThenBB, ElseBB);
1346 CGF.EmitBlock(ThenBB);
1347 // We should reduce with the local Reduce list.
1348 // reduce_function(LocalReduceList, RemoteReduceList)
1349 llvm::Value *LocalDataPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1350 ReduceListAddr.getPointer(), CGF.VoidPtrTy);
1351 llvm::Value *RemoteDataPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1352 RemoteReduceList.getPointer(), CGF.VoidPtrTy);
1353 CGF.EmitCallOrInvoke(ReduceFn, {LocalDataPtr, RemoteDataPtr});
1354 Bld.CreateBr(MergeBB);
1356 CGF.EmitBlock(ElseBB);
1357 // No reduction; just copy:
1358 // Local Reduce list = Remote Reduce list.
1359 emitReductionListCopy(ThreadCopy, CGF, ReductionArrayTy, Privates,
1360 RemoteReduceList, ReduceListAddr);
1361 Bld.CreateBr(MergeBB);
1363 CGF.EmitBlock(MergeBB);
1365 CGF.FinishFunction();
1369 /// This function emits a helper that stores reduced data from the team
1370 /// master to a scratchpad array in global memory.
1372 /// for elem in Reduce List:
1373 /// scratchpad[elem_id][index] = elem
1375 static llvm::Value *emitCopyToScratchpad(CodeGenModule &CGM,
1376 ArrayRef<const Expr *> Privates,
1377 QualType ReductionArrayTy) {
1379 auto &C = CGM.getContext();
1380 auto Int32Ty = C.getIntTypeForBitwidth(32, /* Signed */ true);
1382 // Source of the copy.
1383 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, SourceLocation(),
1384 /*Id=*/nullptr, C.VoidPtrTy);
1385 // Base address of the scratchpad array, with each element storing a
1386 // Reduce list per team.
1387 ImplicitParamDecl ScratchPadArg(C, /*DC=*/nullptr, SourceLocation(),
1388 /*Id=*/nullptr, C.VoidPtrTy);
1389 // A destination index into the scratchpad array, typically the team
1391 ImplicitParamDecl IndexArg(C, /*DC=*/nullptr, SourceLocation(),
1392 /*Id=*/nullptr, Int32Ty);
1393 // Row width of an element in the scratchpad array, typically
1394 // the number of teams.
1395 ImplicitParamDecl WidthArg(C, /*DC=*/nullptr, SourceLocation(),
1396 /*Id=*/nullptr, Int32Ty);
1398 FunctionArgList Args;
1399 Args.push_back(&ReduceListArg);
1400 Args.push_back(&ScratchPadArg);
1401 Args.push_back(&IndexArg);
1402 Args.push_back(&WidthArg);
1404 auto &CGFI = CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
1405 auto *Fn = llvm::Function::Create(
1406 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
1407 "_omp_reduction_copy_to_scratchpad", &CGM.getModule());
1408 CGM.SetInternalFunctionAttributes(/*DC=*/nullptr, Fn, CGFI);
1409 CodeGenFunction CGF(CGM);
1410 // We don't need debug information in this function as nothing here refers to
1412 CGF.disableDebugInfo();
1413 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args);
1415 auto &Bld = CGF.Builder;
1417 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
1418 Address SrcDataAddr(
1419 Bld.CreatePointerBitCastOrAddrSpaceCast(
1420 CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
1421 C.VoidPtrTy, SourceLocation()),
1422 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
1423 CGF.getPointerAlign());
1425 Address AddrScratchPadArg = CGF.GetAddrOfLocalVar(&ScratchPadArg);
1426 llvm::Value *ScratchPadBase = CGF.EmitLoadOfScalar(
1427 AddrScratchPadArg, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
1429 Address AddrIndexArg = CGF.GetAddrOfLocalVar(&IndexArg);
1430 llvm::Value *IndexVal =
1431 Bld.CreateIntCast(CGF.EmitLoadOfScalar(AddrIndexArg, /*Volatile=*/false,
1432 Int32Ty, SourceLocation()),
1433 CGF.SizeTy, /*isSigned=*/true);
1435 Address AddrWidthArg = CGF.GetAddrOfLocalVar(&WidthArg);
1436 llvm::Value *WidthVal =
1437 Bld.CreateIntCast(CGF.EmitLoadOfScalar(AddrWidthArg, /*Volatile=*/false,
1438 Int32Ty, SourceLocation()),
1439 CGF.SizeTy, /*isSigned=*/true);
1441 // The absolute ptr address to the base addr of the next element to copy.
1442 llvm::Value *CumulativeElemBasePtr =
1443 Bld.CreatePtrToInt(ScratchPadBase, CGM.SizeTy);
1444 Address DestDataAddr(CumulativeElemBasePtr, CGF.getPointerAlign());
1446 emitReductionListCopy(ThreadToScratchpad, CGF, ReductionArrayTy, Privates,
1447 SrcDataAddr, DestDataAddr,
1448 {/*RemoteLaneOffset=*/nullptr,
1449 /*ScratchpadIndex=*/IndexVal,
1450 /*ScratchpadWidth=*/WidthVal});
1452 CGF.FinishFunction();
1456 /// This function emits a helper that gathers Reduce lists from the first
1457 /// lane of every active warp to lanes in the first warp.
1459 /// void inter_warp_copy_func(void* reduce_data, num_warps)
1460 /// shared smem[warp_size];
1461 /// For all data entries D in reduce_data:
1462 /// If (I am the first lane in each warp)
1463 /// Copy my local D to smem[warp_id]
1465 /// if (I am the first warp)
1466 /// Copy smem[thread_id] to my local D
1468 static llvm::Value *emitInterWarpCopyFunction(CodeGenModule &CGM,
1469 ArrayRef<const Expr *> Privates,
1470 QualType ReductionArrayTy) {
1471 auto &C = CGM.getContext();
1472 auto &M = CGM.getModule();
1474 // ReduceList: thread local Reduce list.
1475 // At the stage of the computation when this function is called, partially
1476 // aggregated values reside in the first lane of every active warp.
1477 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, SourceLocation(),
1478 /*Id=*/nullptr, C.VoidPtrTy);
1479 // NumWarps: number of warps active in the parallel region. This could
1480 // be smaller than 32 (max warps in a CTA) for partial block reduction.
1481 ImplicitParamDecl NumWarpsArg(C, /*DC=*/nullptr, SourceLocation(),
1483 C.getIntTypeForBitwidth(32, /* Signed */ true));
1484 FunctionArgList Args;
1485 Args.push_back(&ReduceListArg);
1486 Args.push_back(&NumWarpsArg);
1488 auto &CGFI = CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
1489 auto *Fn = llvm::Function::Create(
1490 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
1491 "_omp_reduction_inter_warp_copy_func", &CGM.getModule());
1492 CGM.SetInternalFunctionAttributes(/*DC=*/nullptr, Fn, CGFI);
1493 CodeGenFunction CGF(CGM);
1494 // We don't need debug information in this function as nothing here refers to
1496 CGF.disableDebugInfo();
1497 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args);
1499 auto &Bld = CGF.Builder;
1501 // This array is used as a medium to transfer, one reduce element at a time,
1502 // the data from the first lane of every warp to lanes in the first warp
1503 // in order to perform the final step of a reduction in a parallel region
1504 // (reduction across warps). The array is placed in NVPTX __shared__ memory
1505 // for reduced latency, as well as to have a distinct copy for concurrently
1506 // executing target regions. The array is declared with common linkage so
1507 // as to be shared across compilation units.
1508 const char *TransferMediumName =
1509 "__openmp_nvptx_data_transfer_temporary_storage";
1510 llvm::GlobalVariable *TransferMedium =
1511 M.getGlobalVariable(TransferMediumName);
1512 if (!TransferMedium) {
1513 auto *Ty = llvm::ArrayType::get(CGM.Int64Ty, WarpSize);
1514 unsigned SharedAddressSpace = C.getTargetAddressSpace(LangAS::cuda_shared);
1515 TransferMedium = new llvm::GlobalVariable(
1517 /*isConstant=*/false, llvm::GlobalVariable::CommonLinkage,
1518 llvm::Constant::getNullValue(Ty), TransferMediumName,
1519 /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
1520 SharedAddressSpace);
1523 // Get the CUDA thread id of the current OpenMP thread on the GPU.
1524 auto *ThreadID = getNVPTXThreadID(CGF);
1525 // nvptx_lane_id = nvptx_id % warpsize
1526 auto *LaneID = getNVPTXLaneID(CGF);
1527 // nvptx_warp_id = nvptx_id / warpsize
1528 auto *WarpID = getNVPTXWarpID(CGF);
1530 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
1531 Address LocalReduceList(
1532 Bld.CreatePointerBitCastOrAddrSpaceCast(
1533 CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
1534 C.VoidPtrTy, SourceLocation()),
1535 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
1536 CGF.getPointerAlign());
1539 for (auto &Private : Privates) {
1541 // Warp master copies reduce element to transfer medium in __shared__
1544 llvm::BasicBlock *ThenBB = CGF.createBasicBlock("then");
1545 llvm::BasicBlock *ElseBB = CGF.createBasicBlock("else");
1546 llvm::BasicBlock *MergeBB = CGF.createBasicBlock("ifcont");
1548 // if (lane_id == 0)
1550 Bld.CreateICmpEQ(LaneID, Bld.getInt32(0), "warp_master");
1551 Bld.CreateCondBr(IsWarpMaster, ThenBB, ElseBB);
1552 CGF.EmitBlock(ThenBB);
1554 // Reduce element = LocalReduceList[i]
1555 Address ElemPtrPtrAddr =
1556 Bld.CreateConstArrayGEP(LocalReduceList, Idx, CGF.getPointerSize());
1557 llvm::Value *ElemPtrPtr = CGF.EmitLoadOfScalar(
1558 ElemPtrPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
1559 // elemptr = (type[i]*)(elemptrptr)
1561 Address(ElemPtrPtr, C.getTypeAlignInChars(Private->getType()));
1562 ElemPtr = Bld.CreateElementBitCast(
1563 ElemPtr, CGF.ConvertTypeForMem(Private->getType()));
1565 llvm::Value *Elem = CGF.EmitLoadOfScalar(
1566 ElemPtr, /*Volatile=*/false, Private->getType(), SourceLocation());
1568 // Get pointer to location in transfer medium.
1569 // MediumPtr = &medium[warp_id]
1570 llvm::Value *MediumPtrVal = Bld.CreateInBoundsGEP(
1571 TransferMedium, {llvm::Constant::getNullValue(CGM.Int64Ty), WarpID});
1572 Address MediumPtr(MediumPtrVal, C.getTypeAlignInChars(Private->getType()));
1573 // Casting to actual data type.
1574 // MediumPtr = (type[i]*)MediumPtrAddr;
1575 MediumPtr = Bld.CreateElementBitCast(
1576 MediumPtr, CGF.ConvertTypeForMem(Private->getType()));
1579 Bld.CreateStore(Elem, MediumPtr);
1581 Bld.CreateBr(MergeBB);
1583 CGF.EmitBlock(ElseBB);
1584 Bld.CreateBr(MergeBB);
1586 CGF.EmitBlock(MergeBB);
1588 Address AddrNumWarpsArg = CGF.GetAddrOfLocalVar(&NumWarpsArg);
1589 llvm::Value *NumWarpsVal = CGF.EmitLoadOfScalar(
1590 AddrNumWarpsArg, /*Volatile=*/false, C.IntTy, SourceLocation());
1592 auto *NumActiveThreads = Bld.CreateNSWMul(
1593 NumWarpsVal, getNVPTXWarpSize(CGF), "num_active_threads");
1594 // named_barrier_sync(ParallelBarrierID, num_active_threads)
1595 syncParallelThreads(CGF, NumActiveThreads);
1598 // Warp 0 copies reduce element from transfer medium.
1600 llvm::BasicBlock *W0ThenBB = CGF.createBasicBlock("then");
1601 llvm::BasicBlock *W0ElseBB = CGF.createBasicBlock("else");
1602 llvm::BasicBlock *W0MergeBB = CGF.createBasicBlock("ifcont");
1604 // Up to 32 threads in warp 0 are active.
1605 auto IsActiveThread =
1606 Bld.CreateICmpULT(ThreadID, NumWarpsVal, "is_active_thread");
1607 Bld.CreateCondBr(IsActiveThread, W0ThenBB, W0ElseBB);
1609 CGF.EmitBlock(W0ThenBB);
1611 // SrcMediumPtr = &medium[tid]
1612 llvm::Value *SrcMediumPtrVal = Bld.CreateInBoundsGEP(
1613 TransferMedium, {llvm::Constant::getNullValue(CGM.Int64Ty), ThreadID});
1614 Address SrcMediumPtr(SrcMediumPtrVal,
1615 C.getTypeAlignInChars(Private->getType()));
1616 // SrcMediumVal = *SrcMediumPtr;
1617 SrcMediumPtr = Bld.CreateElementBitCast(
1618 SrcMediumPtr, CGF.ConvertTypeForMem(Private->getType()));
1619 llvm::Value *SrcMediumValue = CGF.EmitLoadOfScalar(
1620 SrcMediumPtr, /*Volatile=*/false, Private->getType(), SourceLocation());
1622 // TargetElemPtr = (type[i]*)(SrcDataAddr[i])
1623 Address TargetElemPtrPtr =
1624 Bld.CreateConstArrayGEP(LocalReduceList, Idx, CGF.getPointerSize());
1625 llvm::Value *TargetElemPtrVal = CGF.EmitLoadOfScalar(
1626 TargetElemPtrPtr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
1627 Address TargetElemPtr =
1628 Address(TargetElemPtrVal, C.getTypeAlignInChars(Private->getType()));
1629 TargetElemPtr = Bld.CreateElementBitCast(
1630 TargetElemPtr, CGF.ConvertTypeForMem(Private->getType()));
1632 // *TargetElemPtr = SrcMediumVal;
1633 CGF.EmitStoreOfScalar(SrcMediumValue, TargetElemPtr, /*Volatile=*/false,
1634 Private->getType());
1635 Bld.CreateBr(W0MergeBB);
1637 CGF.EmitBlock(W0ElseBB);
1638 Bld.CreateBr(W0MergeBB);
1640 CGF.EmitBlock(W0MergeBB);
1642 // While warp 0 copies values from transfer medium, all other warps must
1644 syncParallelThreads(CGF, NumActiveThreads);
1648 CGF.FinishFunction();
1652 /// Emit a helper that reduces data across two OpenMP threads (lanes)
1653 /// in the same warp. It uses shuffle instructions to copy over data from
1654 /// a remote lane's stack. The reduction algorithm performed is specified
1655 /// by the fourth parameter.
1657 /// Algorithm Versions.
1658 /// Full Warp Reduce (argument value 0):
1659 /// This algorithm assumes that all 32 lanes are active and gathers
1660 /// data from these 32 lanes, producing a single resultant value.
1661 /// Contiguous Partial Warp Reduce (argument value 1):
1662 /// This algorithm assumes that only a *contiguous* subset of lanes
1663 /// are active. This happens for the last warp in a parallel region
1664 /// when the user specified num_threads is not an integer multiple of
1665 /// 32. This contiguous subset always starts with the zeroth lane.
1666 /// Partial Warp Reduce (argument value 2):
1667 /// This algorithm gathers data from any number of lanes at any position.
1668 /// All reduced values are stored in the lowest possible lane. The set
1669 /// of problems every algorithm addresses is a super set of those
1670 /// addressable by algorithms with a lower version number. Overhead
1671 /// increases as algorithm version increases.
1675 /// Reduce element refers to the individual data field with primitive
1676 /// data types to be combined and reduced across threads.
1678 /// Reduce list refers to a collection of local, thread-private
1679 /// reduce elements.
1680 /// Remote Reduce list:
1681 /// Remote Reduce list refers to a collection of remote (relative to
1682 /// the current thread) reduce elements.
1684 /// We distinguish between three states of threads that are important to
1685 /// the implementation of this function.
1687 /// Threads in a warp executing the SIMT instruction, as distinguished from
1688 /// threads that are inactive due to divergent control flow.
1690 /// The minimal set of threads that has to be alive upon entry to this
1691 /// function. The computation is correct iff active threads are alive.
1692 /// Some threads are alive but they are not active because they do not
1693 /// contribute to the computation in any useful manner. Turning them off
1694 /// may introduce control flow overheads without any tangible benefits.
1695 /// Effective threads:
1696 /// In order to comply with the argument requirements of the shuffle
1697 /// function, we must keep all lanes holding data alive. But at most
1698 /// half of them perform value aggregation; we refer to this half of
1699 /// threads as effective. The other half is simply handing off their
1704 /// In this step active threads transfer data from higher lane positions
1705 /// in the warp to lower lane positions, creating Remote Reduce list.
1706 /// Value aggregation:
1707 /// In this step, effective threads combine their thread local Reduce list
1708 /// with Remote Reduce list and store the result in the thread local
1711 /// In this step, we deal with the assumption made by algorithm 2
1712 /// (i.e. contiguity assumption). When we have an odd number of lanes
1713 /// active, say 2k+1, only k threads will be effective and therefore k
1714 /// new values will be produced. However, the Reduce list owned by the
1715 /// (2k+1)th thread is ignored in the value aggregation. Therefore
1716 /// we copy the Reduce list from the (2k+1)th lane to (k+1)th lane so
1717 /// that the contiguity assumption still holds.
1718 static llvm::Value *
1719 emitShuffleAndReduceFunction(CodeGenModule &CGM,
1720 ArrayRef<const Expr *> Privates,
1721 QualType ReductionArrayTy, llvm::Value *ReduceFn) {
1722 auto &C = CGM.getContext();
1724 // Thread local Reduce list used to host the values of data to be reduced.
1725 ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, SourceLocation(),
1726 /*Id=*/nullptr, C.VoidPtrTy);
1727 // Current lane id; could be logical.
1728 ImplicitParamDecl LaneIDArg(C, /*DC=*/nullptr, SourceLocation(),
1729 /*Id=*/nullptr, C.ShortTy);
1730 // Offset of the remote source lane relative to the current lane.
1731 ImplicitParamDecl RemoteLaneOffsetArg(C, /*DC=*/nullptr, SourceLocation(),
1732 /*Id=*/nullptr, C.ShortTy);
1733 // Algorithm version. This is expected to be known at compile time.
1734 ImplicitParamDecl AlgoVerArg(C, /*DC=*/nullptr, SourceLocation(),
1735 /*Id=*/nullptr, C.ShortTy);
1736 FunctionArgList Args;
1737 Args.push_back(&ReduceListArg);
1738 Args.push_back(&LaneIDArg);
1739 Args.push_back(&RemoteLaneOffsetArg);
1740 Args.push_back(&AlgoVerArg);
1742 auto &CGFI = CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
1743 auto *Fn = llvm::Function::Create(
1744 CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
1745 "_omp_reduction_shuffle_and_reduce_func", &CGM.getModule());
1746 CGM.SetInternalFunctionAttributes(/*D=*/nullptr, Fn, CGFI);
1747 CodeGenFunction CGF(CGM);
1748 // We don't need debug information in this function as nothing here refers to
1750 CGF.disableDebugInfo();
1751 CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args);
1753 auto &Bld = CGF.Builder;
1755 Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
1756 Address LocalReduceList(
1757 Bld.CreatePointerBitCastOrAddrSpaceCast(
1758 CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
1759 C.VoidPtrTy, SourceLocation()),
1760 CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
1761 CGF.getPointerAlign());
1763 Address AddrLaneIDArg = CGF.GetAddrOfLocalVar(&LaneIDArg);
1764 llvm::Value *LaneIDArgVal = CGF.EmitLoadOfScalar(
1765 AddrLaneIDArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
1767 Address AddrRemoteLaneOffsetArg = CGF.GetAddrOfLocalVar(&RemoteLaneOffsetArg);
1768 llvm::Value *RemoteLaneOffsetArgVal = CGF.EmitLoadOfScalar(
1769 AddrRemoteLaneOffsetArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
1771 Address AddrAlgoVerArg = CGF.GetAddrOfLocalVar(&AlgoVerArg);
1772 llvm::Value *AlgoVerArgVal = CGF.EmitLoadOfScalar(
1773 AddrAlgoVerArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
1775 // Create a local thread-private variable to host the Reduce list
1776 // from a remote lane.
1777 Address RemoteReduceList =
1778 CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.remote_reduce_list");
1780 // This loop iterates through the list of reduce elements and copies,
1781 // element by element, from a remote lane in the warp to RemoteReduceList,
1782 // hosted on the thread's stack.
1783 emitReductionListCopy(RemoteLaneToThread, CGF, ReductionArrayTy, Privates,
1784 LocalReduceList, RemoteReduceList,
1785 {/*RemoteLaneOffset=*/RemoteLaneOffsetArgVal,
1786 /*ScratchpadIndex=*/nullptr,
1787 /*ScratchpadWidth=*/nullptr});
1789 // The actions to be performed on the Remote Reduce list is dependent
1790 // on the algorithm version.
1792 // if (AlgoVer==0) || (AlgoVer==1 && (LaneId < Offset)) || (AlgoVer==2 &&
1793 // LaneId % 2 == 0 && Offset > 0):
1794 // do the reduction value aggregation
1796 // The thread local variable Reduce list is mutated in place to host the
1797 // reduced data, which is the aggregated value produced from local and
1800 // Note that AlgoVer is expected to be a constant integer known at compile
1802 // When AlgoVer==0, the first conjunction evaluates to true, making
1803 // the entire predicate true during compile time.
1804 // When AlgoVer==1, the second conjunction has only the second part to be
1805 // evaluated during runtime. Other conjunctions evaluates to false
1806 // during compile time.
1807 // When AlgoVer==2, the third conjunction has only the second part to be
1808 // evaluated during runtime. Other conjunctions evaluates to false
1809 // during compile time.
1810 auto CondAlgo0 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(0));
1812 auto Algo1 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(1));
1813 auto CondAlgo1 = Bld.CreateAnd(
1814 Algo1, Bld.CreateICmpULT(LaneIDArgVal, RemoteLaneOffsetArgVal));
1816 auto Algo2 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(2));
1817 auto CondAlgo2 = Bld.CreateAnd(
1819 Bld.CreateICmpEQ(Bld.CreateAnd(LaneIDArgVal, Bld.getInt16(1)),
1821 CondAlgo2 = Bld.CreateAnd(
1822 CondAlgo2, Bld.CreateICmpSGT(RemoteLaneOffsetArgVal, Bld.getInt16(0)));
1824 auto CondReduce = Bld.CreateOr(CondAlgo0, CondAlgo1);
1825 CondReduce = Bld.CreateOr(CondReduce, CondAlgo2);
1827 llvm::BasicBlock *ThenBB = CGF.createBasicBlock("then");
1828 llvm::BasicBlock *ElseBB = CGF.createBasicBlock("else");
1829 llvm::BasicBlock *MergeBB = CGF.createBasicBlock("ifcont");
1830 Bld.CreateCondBr(CondReduce, ThenBB, ElseBB);
1832 CGF.EmitBlock(ThenBB);
1833 // reduce_function(LocalReduceList, RemoteReduceList)
1834 llvm::Value *LocalReduceListPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1835 LocalReduceList.getPointer(), CGF.VoidPtrTy);
1836 llvm::Value *RemoteReduceListPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1837 RemoteReduceList.getPointer(), CGF.VoidPtrTy);
1838 CGF.EmitCallOrInvoke(ReduceFn, {LocalReduceListPtr, RemoteReduceListPtr});
1839 Bld.CreateBr(MergeBB);
1841 CGF.EmitBlock(ElseBB);
1842 Bld.CreateBr(MergeBB);
1844 CGF.EmitBlock(MergeBB);
1846 // if (AlgoVer==1 && (LaneId >= Offset)) copy Remote Reduce list to local
1848 Algo1 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(1));
1849 auto CondCopy = Bld.CreateAnd(
1850 Algo1, Bld.CreateICmpUGE(LaneIDArgVal, RemoteLaneOffsetArgVal));
1852 llvm::BasicBlock *CpyThenBB = CGF.createBasicBlock("then");
1853 llvm::BasicBlock *CpyElseBB = CGF.createBasicBlock("else");
1854 llvm::BasicBlock *CpyMergeBB = CGF.createBasicBlock("ifcont");
1855 Bld.CreateCondBr(CondCopy, CpyThenBB, CpyElseBB);
1857 CGF.EmitBlock(CpyThenBB);
1858 emitReductionListCopy(ThreadCopy, CGF, ReductionArrayTy, Privates,
1859 RemoteReduceList, LocalReduceList);
1860 Bld.CreateBr(CpyMergeBB);
1862 CGF.EmitBlock(CpyElseBB);
1863 Bld.CreateBr(CpyMergeBB);
1865 CGF.EmitBlock(CpyMergeBB);
1867 CGF.FinishFunction();
1872 /// Design of OpenMP reductions on the GPU
1874 /// Consider a typical OpenMP program with one or more reduction
1879 /// #pragma omp target teams distribute parallel for \
1880 /// reduction(+:foo) reduction(*:bar)
1881 /// for (int i = 0; i < N; i++) {
1882 /// foo += A[i]; bar *= B[i];
1885 /// where 'foo' and 'bar' are reduced across all OpenMP threads in
1886 /// all teams. In our OpenMP implementation on the NVPTX device an
1887 /// OpenMP team is mapped to a CUDA threadblock and OpenMP threads
1888 /// within a team are mapped to CUDA threads within a threadblock.
1889 /// Our goal is to efficiently aggregate values across all OpenMP
1890 /// threads such that:
1892 /// - the compiler and runtime are logically concise, and
1893 /// - the reduction is performed efficiently in a hierarchical
1894 /// manner as follows: within OpenMP threads in the same warp,
1895 /// across warps in a threadblock, and finally across teams on
1896 /// the NVPTX device.
1898 /// Introduction to Decoupling
1900 /// We would like to decouple the compiler and the runtime so that the
1901 /// latter is ignorant of the reduction variables (number, data types)
1902 /// and the reduction operators. This allows a simpler interface
1903 /// and implementation while still attaining good performance.
1905 /// Pseudocode for the aforementioned OpenMP program generated by the
1906 /// compiler is as follows:
1908 /// 1. Create private copies of reduction variables on each OpenMP
1909 /// thread: 'foo_private', 'bar_private'
1910 /// 2. Each OpenMP thread reduces the chunk of 'A' and 'B' assigned
1911 /// to it and writes the result in 'foo_private' and 'bar_private'
1913 /// 3. Call the OpenMP runtime on the GPU to reduce within a team
1914 /// and store the result on the team master:
1916 /// __kmpc_nvptx_parallel_reduce_nowait(...,
1917 /// reduceData, shuffleReduceFn, interWarpCpyFn)
1920 /// struct ReduceData {
1924 /// reduceData.foo = &foo_private
1925 /// reduceData.bar = &bar_private
1927 /// 'shuffleReduceFn' and 'interWarpCpyFn' are pointers to two
1928 /// auxiliary functions generated by the compiler that operate on
1929 /// variables of type 'ReduceData'. They aid the runtime perform
1930 /// algorithmic steps in a data agnostic manner.
1932 /// 'shuffleReduceFn' is a pointer to a function that reduces data
1933 /// of type 'ReduceData' across two OpenMP threads (lanes) in the
1934 /// same warp. It takes the following arguments as input:
1936 /// a. variable of type 'ReduceData' on the calling lane,
1938 /// c. an offset relative to the current lane_id to generate a
1939 /// remote_lane_id. The remote lane contains the second
1940 /// variable of type 'ReduceData' that is to be reduced.
1941 /// d. an algorithm version parameter determining which reduction
1942 /// algorithm to use.
1944 /// 'shuffleReduceFn' retrieves data from the remote lane using
1945 /// efficient GPU shuffle intrinsics and reduces, using the
1946 /// algorithm specified by the 4th parameter, the two operands
1947 /// element-wise. The result is written to the first operand.
1949 /// Different reduction algorithms are implemented in different
1950 /// runtime functions, all calling 'shuffleReduceFn' to perform
1951 /// the essential reduction step. Therefore, based on the 4th
1952 /// parameter, this function behaves slightly differently to
1953 /// cooperate with the runtime to ensure correctness under
1954 /// different circumstances.
1956 /// 'InterWarpCpyFn' is a pointer to a function that transfers
1957 /// reduced variables across warps. It tunnels, through CUDA
1958 /// shared memory, the thread-private data of type 'ReduceData'
1959 /// from lane 0 of each warp to a lane in the first warp.
1960 /// 4. Call the OpenMP runtime on the GPU to reduce across teams.
1961 /// The last team writes the global reduced value to memory.
1963 /// ret = __kmpc_nvptx_teams_reduce_nowait(...,
1964 /// reduceData, shuffleReduceFn, interWarpCpyFn,
1965 /// scratchpadCopyFn, loadAndReduceFn)
1967 /// 'scratchpadCopyFn' is a helper that stores reduced
1968 /// data from the team master to a scratchpad array in
1971 /// 'loadAndReduceFn' is a helper that loads data from
1972 /// the scratchpad array and reduces it with the input
1975 /// These compiler generated functions hide address
1976 /// calculation and alignment information from the runtime.
1978 /// The team master of the last team stores the reduced
1979 /// result to the globals in memory.
1980 /// foo += reduceData.foo; bar *= reduceData.bar
1983 /// Warp Reduction Algorithms
1985 /// On the warp level, we have three algorithms implemented in the
1986 /// OpenMP runtime depending on the number of active lanes:
1988 /// Full Warp Reduction
1990 /// The reduce algorithm within a warp where all lanes are active
1991 /// is implemented in the runtime as follows:
1993 /// full_warp_reduce(void *reduce_data,
1994 /// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
1995 /// for (int offset = WARPSIZE/2; offset > 0; offset /= 2)
1996 /// ShuffleReduceFn(reduce_data, 0, offset, 0);
1999 /// The algorithm completes in log(2, WARPSIZE) steps.
2001 /// 'ShuffleReduceFn' is used here with lane_id set to 0 because it is
2002 /// not used therefore we save instructions by not retrieving lane_id
2003 /// from the corresponding special registers. The 4th parameter, which
2004 /// represents the version of the algorithm being used, is set to 0 to
2005 /// signify full warp reduction.
2007 /// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
2009 /// #reduce_elem refers to an element in the local lane's data structure
2010 /// #remote_elem is retrieved from a remote lane
2011 /// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
2012 /// reduce_elem = reduce_elem REDUCE_OP remote_elem;
2014 /// Contiguous Partial Warp Reduction
2016 /// This reduce algorithm is used within a warp where only the first
2017 /// 'n' (n <= WARPSIZE) lanes are active. It is typically used when the
2018 /// number of OpenMP threads in a parallel region is not a multiple of
2019 /// WARPSIZE. The algorithm is implemented in the runtime as follows:
2022 /// contiguous_partial_reduce(void *reduce_data,
2023 /// kmp_ShuffleReductFctPtr ShuffleReduceFn,
2024 /// int size, int lane_id) {
2027 /// curr_size = size;
2028 /// mask = curr_size/2;
2029 /// while (offset>0) {
2030 /// ShuffleReduceFn(reduce_data, lane_id, offset, 1);
2031 /// curr_size = (curr_size+1)/2;
2032 /// offset = curr_size/2;
2036 /// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
2038 /// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
2039 /// if (lane_id < offset)
2040 /// reduce_elem = reduce_elem REDUCE_OP remote_elem
2042 /// reduce_elem = remote_elem
2044 /// This algorithm assumes that the data to be reduced are located in a
2045 /// contiguous subset of lanes starting from the first. When there is
2046 /// an odd number of active lanes, the data in the last lane is not
2047 /// aggregated with any other lane's dat but is instead copied over.
2049 /// Dispersed Partial Warp Reduction
2051 /// This algorithm is used within a warp when any discontiguous subset of
2052 /// lanes are active. It is used to implement the reduction operation
2053 /// across lanes in an OpenMP simd region or in a nested parallel region.
2056 /// dispersed_partial_reduce(void *reduce_data,
2057 /// kmp_ShuffleReductFctPtr ShuffleReduceFn) {
2058 /// int size, remote_id;
2059 /// int logical_lane_id = number_of_active_lanes_before_me() * 2;
2061 /// remote_id = next_active_lane_id_right_after_me();
2062 /// # the above function returns 0 of no active lane
2063 /// # is present right after the current lane.
2064 /// size = number_of_active_lanes_in_this_warp();
2065 /// logical_lane_id /= 2;
2066 /// ShuffleReduceFn(reduce_data, logical_lane_id,
2067 /// remote_id-1-threadIdx.x, 2);
2068 /// } while (logical_lane_id % 2 == 0 && size > 1);
2071 /// There is no assumption made about the initial state of the reduction.
2072 /// Any number of lanes (>=1) could be active at any position. The reduction
2073 /// result is returned in the first active lane.
2075 /// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
2077 /// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
2078 /// if (lane_id % 2 == 0 && offset > 0)
2079 /// reduce_elem = reduce_elem REDUCE_OP remote_elem
2081 /// reduce_elem = remote_elem
2084 /// Intra-Team Reduction
2086 /// This function, as implemented in the runtime call
2087 /// '__kmpc_nvptx_parallel_reduce_nowait', aggregates data across OpenMP
2088 /// threads in a team. It first reduces within a warp using the
2089 /// aforementioned algorithms. We then proceed to gather all such
2090 /// reduced values at the first warp.
2092 /// The runtime makes use of the function 'InterWarpCpyFn', which copies
2093 /// data from each of the "warp master" (zeroth lane of each warp, where
2094 /// warp-reduced data is held) to the zeroth warp. This step reduces (in
2095 /// a mathematical sense) the problem of reduction across warp masters in
2096 /// a block to the problem of warp reduction.
2099 /// Inter-Team Reduction
2101 /// Once a team has reduced its data to a single value, it is stored in
2102 /// a global scratchpad array. Since each team has a distinct slot, this
2103 /// can be done without locking.
2105 /// The last team to write to the scratchpad array proceeds to reduce the
2106 /// scratchpad array. One or more workers in the last team use the helper
2107 /// 'loadAndReduceDataFn' to load and reduce values from the array, i.e.,
2108 /// the k'th worker reduces every k'th element.
2110 /// Finally, a call is made to '__kmpc_nvptx_parallel_reduce_nowait' to
2111 /// reduce across workers and compute a globally reduced value.
2113 void CGOpenMPRuntimeNVPTX::emitReduction(
2114 CodeGenFunction &CGF, SourceLocation Loc, ArrayRef<const Expr *> Privates,
2115 ArrayRef<const Expr *> LHSExprs, ArrayRef<const Expr *> RHSExprs,
2116 ArrayRef<const Expr *> ReductionOps, ReductionOptionsTy Options) {
2117 if (!CGF.HaveInsertPoint())
2120 bool ParallelReduction = isOpenMPParallelDirective(Options.ReductionKind);
2121 bool TeamsReduction = isOpenMPTeamsDirective(Options.ReductionKind);
2122 // FIXME: Add support for simd reduction.
2123 assert((TeamsReduction || ParallelReduction) &&
2124 "Invalid reduction selection in emitReduction.");
2126 auto &C = CGM.getContext();
2128 // 1. Build a list of reduction variables.
2129 // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
2130 auto Size = RHSExprs.size();
2131 for (auto *E : Privates) {
2132 if (E->getType()->isVariablyModifiedType())
2133 // Reserve place for array size.
2136 llvm::APInt ArraySize(/*unsigned int numBits=*/32, Size);
2137 QualType ReductionArrayTy =
2138 C.getConstantArrayType(C.VoidPtrTy, ArraySize, ArrayType::Normal,
2139 /*IndexTypeQuals=*/0);
2140 Address ReductionList =
2141 CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.red_list");
2142 auto IPriv = Privates.begin();
2144 for (unsigned I = 0, E = RHSExprs.size(); I < E; ++I, ++IPriv, ++Idx) {
2145 Address Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx,
2146 CGF.getPointerSize());
2147 CGF.Builder.CreateStore(
2148 CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
2149 CGF.EmitLValue(RHSExprs[I]).getPointer(), CGF.VoidPtrTy),
2151 if ((*IPriv)->getType()->isVariablyModifiedType()) {
2152 // Store array size.
2154 Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx,
2155 CGF.getPointerSize());
2156 llvm::Value *Size = CGF.Builder.CreateIntCast(
2158 CGF.getContext().getAsVariableArrayType((*IPriv)->getType()))
2160 CGF.SizeTy, /*isSigned=*/false);
2161 CGF.Builder.CreateStore(CGF.Builder.CreateIntToPtr(Size, CGF.VoidPtrTy),
2166 // 2. Emit reduce_func().
2167 auto *ReductionFn = emitReductionFunction(
2168 CGM, CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo(), Privates,
2169 LHSExprs, RHSExprs, ReductionOps);
2171 // 4. Build res = __kmpc_reduce{_nowait}(<gtid>, <n>, sizeof(RedList),
2172 // RedList, shuffle_reduce_func, interwarp_copy_func);
2173 auto *ThreadId = getThreadID(CGF, Loc);
2174 auto *ReductionArrayTySize = CGF.getTypeSize(ReductionArrayTy);
2175 auto *RL = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
2176 ReductionList.getPointer(), CGF.VoidPtrTy);
2178 auto *ShuffleAndReduceFn = emitShuffleAndReduceFunction(
2179 CGM, Privates, ReductionArrayTy, ReductionFn);
2180 auto *InterWarpCopyFn =
2181 emitInterWarpCopyFunction(CGM, Privates, ReductionArrayTy);
2183 llvm::Value *Res = nullptr;
2184 if (ParallelReduction) {
2185 llvm::Value *Args[] = {ThreadId,
2186 CGF.Builder.getInt32(RHSExprs.size()),
2187 ReductionArrayTySize,
2192 Res = CGF.EmitRuntimeCall(
2193 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_parallel_reduce_nowait),
2197 if (TeamsReduction) {
2198 auto *ScratchPadCopyFn =
2199 emitCopyToScratchpad(CGM, Privates, ReductionArrayTy);
2200 auto *LoadAndReduceFn = emitReduceScratchpadFunction(
2201 CGM, Privates, ReductionArrayTy, ReductionFn);
2203 llvm::Value *Args[] = {ThreadId,
2204 CGF.Builder.getInt32(RHSExprs.size()),
2205 ReductionArrayTySize,
2211 Res = CGF.EmitRuntimeCall(
2212 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_teams_reduce_nowait),
2216 // 5. Build switch(res)
2217 auto *DefaultBB = CGF.createBasicBlock(".omp.reduction.default");
2218 auto *SwInst = CGF.Builder.CreateSwitch(Res, DefaultBB, /*NumCases=*/1);
2220 // 6. Build case 1: where we have reduced values in the master
2221 // thread in each team.
2222 // __kmpc_end_reduce{_nowait}(<gtid>);
2224 auto *Case1BB = CGF.createBasicBlock(".omp.reduction.case1");
2225 SwInst->addCase(CGF.Builder.getInt32(1), Case1BB);
2226 CGF.EmitBlock(Case1BB);
2228 // Add emission of __kmpc_end_reduce{_nowait}(<gtid>);
2229 llvm::Value *EndArgs[] = {ThreadId};
2230 auto &&CodeGen = [&Privates, &LHSExprs, &RHSExprs, &ReductionOps,
2231 this](CodeGenFunction &CGF, PrePostActionTy &Action) {
2232 auto IPriv = Privates.begin();
2233 auto ILHS = LHSExprs.begin();
2234 auto IRHS = RHSExprs.begin();
2235 for (auto *E : ReductionOps) {
2236 emitSingleReductionCombiner(CGF, E, *IPriv, cast<DeclRefExpr>(*ILHS),
2237 cast<DeclRefExpr>(*IRHS));
2243 RegionCodeGenTy RCG(CodeGen);
2244 NVPTXActionTy Action(
2245 nullptr, llvm::None,
2246 createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_end_reduce_nowait),
2248 RCG.setAction(Action);
2250 CGF.EmitBranch(DefaultBB);
2251 CGF.EmitBlock(DefaultBB, /*IsFinished=*/true);