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