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