diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeTRegion.h b/clang/lib/CodeGen/CGOpenMPRuntimeTRegion.h new file mode 100644 --- /dev/null +++ b/clang/lib/CodeGen/CGOpenMPRuntimeTRegion.h @@ -0,0 +1,181 @@ +//===-- CGOpenMPRuntimeTRegion.h --- OpenMP RT TRegion interface codegen --===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Code generation interface for OpenMP target offloading though the generic +// Target Region (TRegion) interface. +// +// See openmp/libomptarget/deviceRTLs/common/target_Region.h for further +// information on the interface functions and their intended use. +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMETREGION_H +#define LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMETREGION_H + +#include "CGOpenMPRuntimeTarget.h" +#include "llvm/ADT/SmallBitVector.h" + +namespace clang { +namespace CodeGen { + +class CGOpenMPRuntimeTRegion : public CGOpenMPRuntimeTarget { + // TODO: The target region interface only covers kernel codes for now. This + // therefore codegen implicitly assumes the target region kernel + // interface is targeted. Once a second target region interface is put + // in place, e.g., specialized to many-core offloading, we might need + // to make the target interface explicit. + + /// Create an outlined function for a target kernel. + /// + /// \param D Directive to emit. + /// \param ParentName Name of the function that encloses the target region. + /// \param OutlinedFn Outlined function value to be defined by this call. + /// \param OutlinedFnID Outlined function ID value to be defined by this call. + /// \param CodeGen Object containing the target statements. + /// An outlined function may not be an entry if, e.g. the if clause always + /// evaluates to false. + void emitKernel(const OMPExecutableDirective &D, StringRef ParentName, + llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID, + const RegionCodeGenTy &CodeGen); + + /// Helper for generic kernel mode, target directive's entry function. + void emitKernelHeader(CodeGenFunction &CGF, llvm::BasicBlock *&ExitBB); + + /// Signal termination of generic mode execution. + void emitKernelFooter(CodeGenFunction &CGF, llvm::BasicBlock *ExitBB); + + // + // Base class overrides. + // + + /// Emit outlined function for 'target' directive. + /// + /// \param D Directive to emit. + /// \param ParentName Name of the function that encloses the target region. + /// \param OutlinedFn Outlined function value to be defined by this call. + /// \param OutlinedFnID Outlined function ID value to be defined by this call. + /// \param IsOffloadEntry True if the outlined function is an offload entry. + /// An outlined function may not be an entry if, e.g. the if clause always + /// evaluates to false. + void emitTargetOutlinedFunction(const OMPExecutableDirective &D, + StringRef ParentName, + llvm::Function *&OutlinedFn, + llvm::Constant *&OutlinedFnID, + bool IsOffloadEntry, + const RegionCodeGenTy &CodeGen) override; + +protected: + /// Get the function name of an outlined region, customized to the target. + StringRef getOutlinedHelperName() const override { return ".omp_TRegion."; } + +public: + explicit CGOpenMPRuntimeTRegion(CodeGenModule &CGM); + + /// Emits inlined function for the specified OpenMP parallel directive. + /// + /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID, + /// kmp_int32 BoundID, struct context_vars*). + /// \param D OpenMP directive. + /// \param ThreadIDVar Variable for thread id in the current OpenMP region. + /// \param InnermostKind Kind of innermost directive (for simple directives it + /// is a directive itself, for combined - its innermost directive). + /// \param CodeGen Code generation sequence for the \a D directive. + llvm::Function * + emitParallelOutlinedFunction(const OMPExecutableDirective &D, + const VarDecl *ThreadIDVar, + OpenMPDirectiveKind InnermostKind, + const RegionCodeGenTy &CodeGen) override; + + /// Emits code for parallel or serial call of the \a OutlinedFn with + /// variables captured in a record which address is stored in \a + /// CapturedStruct. + /// \param OutlinedFn Outlined function to be run in parallel threads. Type of + /// this function is void(*)(kmp_int32 *, kmp_int32, struct context_vars*). + /// \param CapturedVars A pointer to the record with the references to + /// variables used in \a OutlinedFn function. + /// \param IfCond Condition in the associated 'if' clause, if it was + /// specified, nullptr otherwise. + void emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc, + llvm::Function *OutlinedFn, + ArrayRef CapturedVars, + const Expr *IfCond) override; + + /// Emits a critical region. + /// \param CriticalName Name of the critical region. + /// \param CriticalOpGen Generator for the statement associated with the given + /// critical region. + /// \param Hint Value of the 'hint' clause (optional). + void emitCriticalRegion(CodeGenFunction &CGF, StringRef CriticalName, + const RegionCodeGenTy &CriticalOpGen, + SourceLocation Loc, + const Expr *Hint = nullptr) override; + + /// Emit a code for reduction clause. + /// + /// \param Privates List of private copies for original reduction arguments. + /// \param LHSExprs List of LHS in \a ReductionOps reduction operations. + /// \param RHSExprs List of RHS in \a ReductionOps reduction operations. + /// \param ReductionOps List of reduction operations in form 'LHS binop RHS' + /// or 'operator binop(LHS, RHS)'. + /// \param Options List of options for reduction codegen: + /// WithNowait true if parent directive has also nowait clause, false + /// otherwise. + /// SimpleReduction Emit reduction operation only. Used for omp simd + /// directive on the host. + /// ReductionKind The kind of reduction to perform. + virtual void emitReduction(CodeGenFunction &CGF, SourceLocation Loc, + ArrayRef Privates, + ArrayRef LHSExprs, + ArrayRef RHSExprs, + ArrayRef ReductionOps, + ReductionOptionsTy Options) override; + +protected: + + /// Hook to allow derived classes to perform checks on the AST that justify + /// execution without runtime support. + virtual bool mayNeedRuntimeSupport() const { return true; } + + /// Hook to allow derived classes to perform checks on the AST that justify + /// execution without data sharing support. + virtual bool mayPerformDataSharing() const { return true; } + +private: + + /// Helper to check if SPMD mode is enabled. Derived classes that perform + /// checks on the AST to justify SPMD mode can overload the + /// CGOpenMPRuntimeTarget::getExecutionMode(). + bool isKnownSPMDMode() const { return getExecutionMode() == EM_SPMD; } + + /// Simple container for a wrapper of an outlined parallel function and the + /// layout of the passed variables (= captured variables, both shared and + /// firstprivate). + struct WrapperInfo { + llvm::Function *WrapperFn = nullptr; + llvm::StructType *SharedVarsStructTy = nullptr; + llvm::StructType *PrivateVarsStructTy = nullptr; + llvm::SmallBitVector CaptureIsPrivate; + }; + + /// Map an outlined function to its wrapper and shared struct type. The latter + /// defines the layout of the payload and the wrapper will unpack that payload + /// and pass the values to the outlined function. + llvm::DenseMap WrapperInfoMap; + + /// Emit function which wraps the outline parallel region + /// and controls the parameters which are passed to this function. + /// The wrapper ensures that the outlined function is called + /// with the correct arguments when data is shared. + void createParallelDataSharingWrapper(llvm::Function *OutlinedParallelFn, + const OMPExecutableDirective &D); +}; + +} // namespace CodeGen +} // namespace clang + +#endif // LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMEKERNEL_H diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeTRegion.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeTRegion.cpp new file mode 100644 --- /dev/null +++ b/clang/lib/CodeGen/CGOpenMPRuntimeTRegion.cpp @@ -0,0 +1,441 @@ +//===-- CGOpenMPRuntimeTRegion.cpp - OpenMP RT TRegion interface codegen --===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Implementation of the code generation interface for OpenMP target offloading +// though the Target Region (TRegion) interface. +// +// See the file comment in CGOpenMPRuntimeTRegion.h for more information. +// +//===----------------------------------------------------------------------===// + +#include "CGOpenMPRuntimeTRegion.h" +#include "CodeGenFunction.h" +#include "clang/AST/StmtVisitor.h" + +using namespace clang; +using namespace CodeGen; + +void CGOpenMPRuntimeTRegion::emitKernel(const OMPExecutableDirective &D, + StringRef ParentName, + llvm::Function *&OutlinedFn, + llvm::Constant *&OutlinedFnID, + const RegionCodeGenTy &CodeGen) { + WrapperInfoMap.clear(); + + // Emit target region as a standalone region. + class KernelPrePostActionTy : public PrePostActionTy { + CGOpenMPRuntimeTRegion &RT; + llvm::BasicBlock *ExitBB; + + public: + KernelPrePostActionTy(CGOpenMPRuntimeTRegion &RT) + : RT(RT), ExitBB(nullptr) {} + + void Enter(CodeGenFunction &CGF) override { + RT.emitKernelHeader(CGF, ExitBB); + // Skip target region initialization. + RT.setLocThreadIdInsertPt(CGF, /* AtCurrentPoint */ true); + } + + void Exit(CodeGenFunction &CGF) override { + RT.clearLocThreadIdInsertPt(CGF); + RT.emitKernelFooter(CGF, ExitBB); + } + + } Action(*this); + CodeGen.setAction(Action); + + emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID, + /* IsOffloadEntry */ true, CodeGen); +} + +void CGOpenMPRuntimeTRegion::emitKernelHeader(CodeGenFunction &CGF, + llvm::BasicBlock *&ExitBB) { + CGBuilderTy &Bld = CGF.Builder; + + // Setup BBs in entry function. + llvm::BasicBlock *ExecuteBB = CGF.createBasicBlock(".execute"); + ExitBB = CGF.createBasicBlock(".exit"); + + llvm::Value *Args[] = { + /* Ident */ llvm::Constant::getNullValue(getIdentTyPointerTy()), + /* UseSPMDMode */ Bld.getInt1(isKnownSPMDMode()), + /* UseStateMachine */ Bld.getInt1(1), + /* RequiresOMPRuntime */ + Bld.getInt1(mayNeedRuntimeSupport()), + /* RequiresDataSharing */ Bld.getInt1(mayPerformDataSharing())}; + llvm::CallInst *InitCI = CGF.EmitRuntimeCall( + createTargetRuntimeFunction(OMPRTL__kmpc_target_region_kernel_init), + Args); + + llvm::Value *ExecuteCnd = Bld.CreateICmpEQ(InitCI, Bld.getInt8(1)); + + Bld.CreateCondBr(ExecuteCnd, ExecuteBB, ExitBB); + CGF.EmitBlock(ExecuteBB); +} + +void CGOpenMPRuntimeTRegion::emitKernelFooter(CodeGenFunction &CGF, + llvm::BasicBlock *ExitBB) { + if (!CGF.HaveInsertPoint()) + return; + + llvm::BasicBlock *OMPDeInitBB = CGF.createBasicBlock(".omp.deinit"); + CGF.EmitBranch(OMPDeInitBB); + + CGF.EmitBlock(OMPDeInitBB); + + CGBuilderTy &Bld = CGF.Builder; + // DeInitialize the OMP state in the runtime; called by all active threads. + llvm::Value *Args[] = { + /* Ident */ llvm::Constant::getNullValue(getIdentTyPointerTy()), + /* UseSPMDMode */ Bld.getInt1(isKnownSPMDMode()), + /* RequiredOMPRuntime */ + Bld.getInt1(mayNeedRuntimeSupport())}; + + CGF.EmitRuntimeCall( + createTargetRuntimeFunction(OMPRTL__kmpc_target_region_kernel_deinit), + Args); + + CGF.EmitBranch(ExitBB); + CGF.EmitBlock(ExitBB); +} + +void CGOpenMPRuntimeTRegion::emitTargetOutlinedFunction( + const OMPExecutableDirective &D, StringRef ParentName, + llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID, + bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) { + if (!IsOffloadEntry) // Nothing to do. + return; + + assert(!ParentName.empty() && "Invalid target region parent name!"); + + emitKernel(D, ParentName, OutlinedFn, OutlinedFnID, CodeGen); + + // Create a unique global variable to indicate the execution mode of this + // target region. The execution mode is either 'non-SPMD' or 'SPMD'. Initially + // all regions are executed in non-SPMD mode. This variable is picked up by + // the offload library to setup the device appropriately before kernel launch. + auto *GVMode = new llvm::GlobalVariable( + CGM.getModule(), CGM.Int8Ty, /* isConstant */ true, + llvm::GlobalValue::WeakAnyLinkage, llvm::ConstantInt::get(CGM.Int8Ty, 1), + Twine(OutlinedFn->getName(), "_exec_mode")); + CGM.addCompilerUsedGlobal(GVMode); +} + +CGOpenMPRuntimeTRegion::CGOpenMPRuntimeTRegion(CodeGenModule &CGM) + : CGOpenMPRuntimeTarget(CGM) { + if (!CGM.getLangOpts().OpenMPIsDevice) + llvm_unreachable("TRegion code generation does only handle device code!"); +} + +llvm::Function *CGOpenMPRuntimeTRegion::emitParallelOutlinedFunction( + const OMPExecutableDirective &D, const VarDecl *ThreadIDVar, + OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) { + + // Emit target region as a standalone region. + llvm::Function *OutlinedFun = + cast(CGOpenMPRuntime::emitParallelOutlinedFunction( + D, ThreadIDVar, InnermostKind, CodeGen)); + + createParallelDataSharingWrapper(OutlinedFun, D); + + return OutlinedFun; +} + +void CGOpenMPRuntimeTRegion::createParallelDataSharingWrapper( + llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) { + ASTContext &Ctx = CGM.getContext(); + const auto &CS = *D.getCapturedStmt(OMPD_parallel); + + // Create a function that takes as argument the source thread. + FunctionArgList WrapperArgs; + ImplicitParamDecl SharedVarsArgDecl(Ctx, /* DC */ nullptr, D.getBeginLoc(), + /* Id */ nullptr, Ctx.VoidPtrTy, + ImplicitParamDecl::Other); + ImplicitParamDecl PrivateVarsArgDecl(Ctx, /* DC */ nullptr, D.getBeginLoc(), + /* Id */ nullptr, Ctx.VoidPtrTy, + ImplicitParamDecl::Other); + WrapperArgs.emplace_back(&SharedVarsArgDecl); + WrapperArgs.emplace_back(&PrivateVarsArgDecl); + + const CGFunctionInfo &CGFI = + CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, WrapperArgs); + + auto *WrapperFn = llvm::Function::Create( + CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage, + Twine(OutlinedParallelFn->getName(), "_wrapper"), &CGM.getModule()); + CGM.SetInternalFunctionAttributes(GlobalDecl(), WrapperFn, CGFI); + + OutlinedParallelFn->setLinkage(llvm::GlobalValue::InternalLinkage); + OutlinedParallelFn->setDoesNotRecurse(); + WrapperFn->setLinkage(llvm::GlobalValue::InternalLinkage); + WrapperFn->setDoesNotRecurse(); + + CodeGenFunction CGF(CGM, /* suppressNewContext */ true); + CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, WrapperFn, CGFI, WrapperArgs, + D.getBeginLoc(), D.getBeginLoc()); + + auto AI = WrapperFn->arg_begin(); + llvm::Argument &SharedVarsArg = *(AI++); + llvm::Argument &PrivateVarsArg = *(AI); + SharedVarsArg.setName("shared_vars"); + PrivateVarsArg.setName("private_vars"); + + setLocThreadIdInsertPt(CGF, /* AtCurrentPoint */ true); + + Address ThreadIDAddr = emitThreadIDAddress(CGF, D.getBeginLoc()); + Address ZeroAddr = CGF.CreateMemTemp(CGF.getContext().getIntTypeForBitwidth( + /* DestWidth */ 32, /* Signed */ 1), + /* Name */ ".zero.addr"); + CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/* C */ 0)); + + // Create the array of arguments and fill it with boilerplate values. + SmallVector Args; + Args.emplace_back(ThreadIDAddr.getPointer()); + Args.emplace_back(ZeroAddr.getPointer()); + + CGBuilderTy &Bld = CGF.Builder; + + // Collect all variables marked as shared. + llvm::SmallPtrSet SharedVars; + for (const auto *C : D.getClausesOfKind()) + for (const Expr *E : C->getVarRefs()) + SharedVars.insert(CGOpenMPRuntimeTarget::getUnderlyingVar(E)); + + // Retrieve the shared and private variables from argument pointers and pass + // them to the outlined function. + llvm::SmallVector SharedStructMemberTypes; + llvm::SmallVector PrivateStructMemberTypes; + + WrapperInfo &WI = WrapperInfoMap[OutlinedParallelFn]; + WI.WrapperFn = WrapperFn; + + llvm::LLVMContext &LLVMCtx = OutlinedParallelFn->getContext(); + auto ArgIt = OutlinedParallelFn->arg_begin(); + ArgIt->addAttr(llvm::Attribute::get( + LLVMCtx, llvm::Attribute::Dereferenceable, + ThreadIDAddr.getElementType()->getScalarSizeInBits() / 8)); + ArgIt->addAttr( + llvm::Attribute::get(LLVMCtx, llvm::Attribute::Alignment, + ThreadIDAddr.getAlignment().getQuantity())); + ArgIt++; + + ArgIt->addAttr(llvm::Attribute::get( + LLVMCtx, llvm::Attribute::Dereferenceable, + ZeroAddr.getElementType()->getScalarSizeInBits() / 8)); + ArgIt->addAttr(llvm::Attribute::get(LLVMCtx, llvm::Attribute::Alignment, + ZeroAddr.getAlignment().getQuantity())); + ArgIt++; + + // If we require loop bounds they are already part of the outlined function + // encoding, just after global_tid and bound_tid. + bool RequiresLoopBounds = + isOpenMPLoopBoundSharingDirective(D.getDirectiveKind()); + if (RequiresLoopBounds) { + // Register the lower bound in the wrapper info. + WI.CaptureIsPrivate.push_back(true); + PrivateStructMemberTypes.push_back((ArgIt++)->getType()); + // Register the upper bound in the wrapper info. + WI.CaptureIsPrivate.push_back(true); + PrivateStructMemberTypes.push_back((ArgIt++)->getType()); + } + + auto CIt = CS.capture_begin(); + for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CIt) { + bool IsPrivate = CIt->capturesVariableArrayType() || + CIt->capturesVariableByCopy() || + !SharedVars.count(CIt->getCapturedVar()); + WI.CaptureIsPrivate.push_back(IsPrivate); + + auto &StructMemberTypes = + IsPrivate ? PrivateStructMemberTypes : SharedStructMemberTypes; + llvm::Type *ArgTy = (ArgIt++)->getType(); + if (!IsPrivate) { + assert(ArgTy->isPointerTy()); + ArgTy = ArgTy->getPointerElementType(); + } + StructMemberTypes.push_back(ArgTy); + } + + // Verify the position of the outlined function argument iterator as a sanity + // check. + assert(ArgIt == OutlinedParallelFn->arg_end() && + "Not all arguments have been processed!"); + + llvm::Value *SharedVarsStructPtr = nullptr; + llvm::Value *PrivateVarsStructPtr = nullptr; + if (!PrivateStructMemberTypes.empty()) { + WI.PrivateVarsStructTy = llvm::StructType::create( + LLVMCtx, PrivateStructMemberTypes, "omp.private.struct"); + PrivateVarsStructPtr = Bld.CreateBitCast( + &PrivateVarsArg, WI.PrivateVarsStructTy->getPointerTo()); + } + if (!SharedStructMemberTypes.empty()) { + WI.SharedVarsStructTy = llvm::StructType::create( + LLVMCtx, SharedStructMemberTypes, "omp.shared.struct"); + SharedVarsStructPtr = Bld.CreateBitCast( + &SharedVarsArg, WI.SharedVarsStructTy->getPointerTo()); + } + + assert(WI.CaptureIsPrivate.size() + /* global_tid & bound_tid */ 2 == + OutlinedParallelFn->arg_size() && + "Not all arguments have been processed!"); + + unsigned PrivateIdx = 0, SharedIdx = 0; + for (int i = 0, e = WI.CaptureIsPrivate.size(); i < e; i++) { + bool IsPrivate = WI.CaptureIsPrivate[i]; + + llvm::Value *StructPtr = + IsPrivate ? PrivateVarsStructPtr : SharedVarsStructPtr; + unsigned &Idx = IsPrivate ? PrivateIdx : SharedIdx; + + // TODO: Figure out the real alignment + if (IsPrivate) { + Args.emplace_back( + Bld.CreateAlignedLoad(Bld.CreateStructGEP(StructPtr, Idx++), 1)); + } else { + llvm::Value *GEP = Bld.CreateStructGEP(StructPtr, Idx++); + Args.emplace_back(GEP); + } + } + + assert(Args.size() == OutlinedParallelFn->arg_size()); + emitOutlinedFunctionCall(CGF, D.getBeginLoc(), OutlinedParallelFn, Args); + + CGF.FinishFunction(); + + clearLocThreadIdInsertPt(CGF); +} + +void CGOpenMPRuntimeTRegion::emitParallelCall( + CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *Fn, + ArrayRef CapturedVars, const Expr *IfCond) { + if (!CGF.HaveInsertPoint()) + return; + + const WrapperInfo &WI = WrapperInfoMap[Fn]; + + auto &&ParGen = [this, CapturedVars, WI](CodeGenFunction &CGF, + PrePostActionTy &) { + CGBuilderTy &Bld = CGF.Builder; + assert(WI.WrapperFn && "Wrapper function does not exist!"); + + llvm::Value *SharedVarsSize = llvm::Constant::getNullValue(CGM.Int16Ty); + llvm::Value *PrivateVarsSize = SharedVarsSize; + llvm::Value *SharedStructAlloca = llvm::UndefValue::get(CGM.VoidPtrTy); + llvm::Value *PrivateStructAlloca = SharedStructAlloca; + + if (WI.SharedVarsStructTy) { + SharedStructAlloca = CGF.CreateDefaultAlignTempAlloca( + WI.SharedVarsStructTy, ".shared.vars") + .getPointer(); + const llvm::DataLayout &DL = WI.WrapperFn->getParent()->getDataLayout(); + SharedVarsSize = Bld.getInt16(DL.getTypeAllocSize(WI.SharedVarsStructTy)); + } + if (WI.PrivateVarsStructTy) { + PrivateStructAlloca = CGF.CreateDefaultAlignTempAlloca( + WI.PrivateVarsStructTy, ".private.vars") + .getPointer(); + const llvm::DataLayout &DL = WI.WrapperFn->getParent()->getDataLayout(); + PrivateVarsSize = + Bld.getInt16(DL.getTypeAllocSize(WI.PrivateVarsStructTy)); + } + + llvm::SmallVector Args; + Args.push_back( + /* Ident */ llvm::Constant::getNullValue(getIdentTyPointerTy())); + Args.push_back( + /* UseSPMDMode */ Bld.getInt16(getExecutionMode())); + Args.push_back( + /* RequiredOMPRuntime */ Bld.getInt1(mayNeedRuntimeSupport())); + Args.push_back(WI.WrapperFn); + Args.push_back(CGF.EmitCastToVoidPtr(SharedStructAlloca)); + Args.push_back(SharedVarsSize); + Args.push_back(CGF.EmitCastToVoidPtr(PrivateStructAlloca)); + Args.push_back(PrivateVarsSize); + Args.push_back( + /* SharedPointers */ Bld.getInt1(0)); + + assert((CapturedVars.empty() || + (WI.SharedVarsStructTy || WI.PrivateVarsStructTy)) && + "Expected the shared or private struct type to be set if variables " + "are captured!"); + assert((CapturedVars.empty() || + CapturedVars.size() == + (WI.SharedVarsStructTy ? WI.SharedVarsStructTy->getNumElements() + : 0) + + (WI.PrivateVarsStructTy + ? WI.PrivateVarsStructTy->getNumElements() + : 0)) && + "# elements in shared struct types should be number of captured " + "variables!"); + + // Store all captured variables into a single local structure that is then + // passed to the runtime library. + unsigned PrivateIdx = 0, SharedIdx = 0; + for (int i = 0, e = WI.CaptureIsPrivate.size(); i < e; i++) { + bool IsPrivate = WI.CaptureIsPrivate[i]; + + llvm::Value *StructPtr = + IsPrivate ? PrivateStructAlloca : SharedStructAlloca; + unsigned &Idx = IsPrivate ? PrivateIdx : SharedIdx; + llvm::Value *GEP = Bld.CreateStructGEP(StructPtr, Idx++); + llvm::Value *Var = IsPrivate ? CapturedVars[i] + : Bld.CreateAlignedLoad(CapturedVars[i], 1); + Bld.CreateDefaultAlignedStore(Var, GEP); + } + + CGF.EmitRuntimeCall( + createTargetRuntimeFunction(OMPRTL__kmpc_target_region_kernel_parallel), + Args); + + SharedIdx = 0; + for (int i = 0, e = WI.CaptureIsPrivate.size(); i < e; i++) { + bool IsPrivate = WI.CaptureIsPrivate[i]; + if (IsPrivate) + continue; + + llvm::Value *GEP = Bld.CreateStructGEP(SharedStructAlloca, SharedIdx++); + llvm::Value *Var = Bld.CreateAlignedLoad(GEP, 1); + Bld.CreateDefaultAlignedStore(Var, CapturedVars[i]); + } + }; + + auto &&SeqGen = [this, &ParGen, Loc](CodeGenFunction &CGF, + PrePostActionTy &Action) { + // Use an artifical "num_threads(1)" clause to force sequential execution if + // the expression in the 'if clause' evaluated to false. We expect the + // middle-end to clean this up. + emitNumThreadsClause(CGF, CGF.Builder.getInt32(/* C */ 1), Loc); + ParGen(CGF, Action); + }; + + if (IfCond) { + emitOMPIfClause(CGF, IfCond, ParGen, SeqGen); + } else { + CodeGenFunction::RunCleanupsScope Scope(CGF); + RegionCodeGenTy ThenRCG(ParGen); + ThenRCG(CGF); + } +} + +void CGOpenMPRuntimeTRegion::emitCriticalRegion( + CodeGenFunction &CGF, StringRef CriticalName, + const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc, + const Expr *Hint) { + llvm_unreachable( + "TODO: TRegion code generation does not support critical regions yet!"); +} + +void CGOpenMPRuntimeTRegion::emitReduction( + CodeGenFunction &CGF, SourceLocation Loc, ArrayRef Privates, + ArrayRef LHSExprs, ArrayRef RHSExprs, + ArrayRef ReductionOps, ReductionOptionsTy Options) { + llvm_unreachable( + "TODO: TRegion code generation does not support reductions yet!"); +} diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeTarget.h b/clang/lib/CodeGen/CGOpenMPRuntimeTarget.h --- a/clang/lib/CodeGen/CGOpenMPRuntimeTarget.h +++ b/clang/lib/CodeGen/CGOpenMPRuntimeTarget.h @@ -25,11 +25,11 @@ /// Defines the execution mode. enum ExecutionMode { /// SPMD execution mode (all threads are worker threads). - EM_SPMD, + EM_SPMD = 1, /// Non-SPMD execution mode (1 master thread, others are workers). - EM_NonSPMD, + EM_NonSPMD = 0, /// Unknown execution mode (orphaned directive). - EM_Unknown, + EM_Unknown = -1, }; /// Return the execution mode, if not overloaded this is always Unknown. diff --git a/clang/lib/CodeGen/CMakeLists.txt b/clang/lib/CodeGen/CMakeLists.txt --- a/clang/lib/CodeGen/CMakeLists.txt +++ b/clang/lib/CodeGen/CMakeLists.txt @@ -70,6 +70,7 @@ CGOpenMPRuntime.cpp CGOpenMPRuntimeNVPTX.cpp CGOpenMPRuntimeTarget.cpp + CGOpenMPRuntimeTRegion.cpp CGRecordLayoutBuilder.cpp CGStmt.cpp CGStmtOpenMP.cpp diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -20,6 +20,7 @@ #include "CGOpenCLRuntime.h" #include "CGOpenMPRuntime.h" #include "CGOpenMPRuntimeNVPTX.h" +#include "CGOpenMPRuntimeTRegion.h" #include "CodeGenFunction.h" #include "CodeGenPGO.h" #include "ConstantEmitter.h" @@ -67,6 +68,11 @@ llvm::cl::desc("Emit limited coverage mapping information (experimental)"), llvm::cl::init(false)); +static llvm::cl::opt UseGenericTRegionInterface( + "openmp-tregion-runtime", llvm::cl::ZeroOrMore, llvm::cl::Hidden, + llvm::cl::desc("Use the generic target region OpenMP runtime interface"), + llvm::cl::init(false)); + static const char AnnotationSection[] = "llvm.metadata"; static CGCXXABI *createCXXABI(CodeGenModule &CGM) { @@ -206,7 +212,10 @@ case llvm::Triple::nvptx64: assert(getLangOpts().OpenMPIsDevice && "OpenMP NVPTX is only prepared to deal with device code."); - OpenMPRuntime.reset(new CGOpenMPRuntimeNVPTX(*this)); + if (UseGenericTRegionInterface) + OpenMPRuntime.reset(new CGOpenMPRuntimeTRegion(*this)); + else + OpenMPRuntime.reset(new CGOpenMPRuntimeNVPTX(*this)); break; default: if (LangOpts.OpenMPSimd) diff --git a/clang/test/OpenMP/target_tregion_no_SPMD_mode.c b/clang/test/OpenMP/target_tregion_no_SPMD_mode.c new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/target_tregion_no_SPMD_mode.c @@ -0,0 +1,72 @@ +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -fopenmp -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -mllvm -openmp-tregion-runtime -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s + +// CHECK: loop_in_loop_in_tregion +// CHECK: %0 = call i8 @__kmpc_target_region_kernel_init(%struct.ident_t* null, i1 false, i1 true, i1 true, i1 true) +// CHECK: call void @__kmpc_target_region_kernel_deinit(%struct.ident_t* null, i1 false, i1 true) +void loop_in_loop_in_tregion(int *A, int *B) { +#pragma omp target + for (int i = 0; i < 512; i++) { + for (int j = 0; j < 1024; j++) + A[j] += B[i + j]; + } +} + +// CHECK: parallel_loops_and_accesses_in_tregion +// CHECK: %0 = call i8 @__kmpc_target_region_kernel_init(%struct.ident_t* null, i1 false, i1 true, i1 true, i1 true) +// CHECK: call void @__kmpc_target_region_kernel_parallel(%struct.ident_t* null, i1 false, i1 true, void (i8*, i8*)* @.omp_TRegion._wrapper, i8* undef, i16 0, i8* %2, i16 16, i1 false) +// CHECK: call void @__kmpc_target_region_kernel_parallel(%struct.ident_t* null, i1 false, i1 true, void (i8*, i8*)* @.omp_TRegion.1_wrapper, i8* undef, i16 0, i8* %5, i16 16, i1 false) +// CHECK: call void @__kmpc_target_region_kernel_parallel(%struct.ident_t* null, i1 false, i1 true, void (i8*, i8*)* @.omp_TRegion.2_wrapper, i8* undef, i16 0, i8* %8, i16 16, i1 false) +// CHECK: call void @__kmpc_target_region_kernel_deinit(%struct.ident_t* null, i1 false, i1 true) +void parallel_loops_and_accesses_in_tregion(int *A, int *B) { +#pragma omp target + { +#pragma omp parallel for + for (int j = 0; j < 1024; j++) + A[j] += B[0 + j]; +#pragma omp parallel for + for (int j = 0; j < 1024; j++) + A[j] += B[1 + j]; +#pragma omp parallel for + for (int j = 0; j < 1024; j++) + A[j] += B[2 + j]; + + // This needs a guard in SPMD mode + A[0] = B[0]; + } +} + +void extern_func(); +static void parallel_loop(int *A, int *B, int i) { +#pragma omp parallel for + for (int j = 0; j < 1024; j++) + A[j] += B[i + j]; +} + +// CHECK: parallel_loop_in_function_in_loop_with_global_acc_in_tregion +// CHECK: %1 = call i8 @__kmpc_target_region_kernel_init(%struct.ident_t* null, i1 false, i1 true, i1 true, i1 true) +// CHECK: call void @__kmpc_target_region_kernel_deinit(%struct.ident_t* null, i1 false, i1 true) +int Global[512]; +void parallel_loop_in_function_in_loop_with_global_acc_in_tregion(int *A, int *B) { +#pragma omp target + for (int i = 0; i < 512; i++) { + parallel_loop(A, B, i); + Global[i]++; + } +} + +// CHECK: parallel_loop +// CHECK: call void @__kmpc_target_region_kernel_parallel(%struct.ident_t* null, i1 false, i1 true, void (i8*, i8*)* @.omp_TRegion.3_wrapper, i8* undef, i16 0, i8* %0, i16 24, i1 false) + +// CHECK: parallel_loops_in_functions_and_extern_func_in_tregion +// CHECK: %0 = call i8 @__kmpc_target_region_kernel_init(%struct.ident_t* null, i1 false, i1 true, i1 true, i1 true) +// CHECK: call void @__kmpc_target_region_kernel_deinit(%struct.ident_t* null, i1 false, i1 true) +void parallel_loops_in_functions_and_extern_func_in_tregion(int *A, int *B) { +#pragma omp target + { + parallel_loop(A, B, 1); + parallel_loop(A, B, 2); + extern_func(); + parallel_loop(A, B, 3); + } +}