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