From: Alexey Bataev Date: Wed, 25 May 2016 12:36:08 +0000 (+0000) Subject: [OPENMP 4.5] Codegen for dacross loop synchronization constructs. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=f2b5110519dc87a957877b8bdaf26717ff25abb5;p=clang [OPENMP 4.5] Codegen for dacross loop synchronization constructs. OpenMP 4.5 adds support for doacross loop synchronization. Patch implements codegen for this construct. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@270690 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/AST/OpenMPClause.h b/include/clang/AST/OpenMPClause.h index c11a4d6e1d..13c05ca634 100644 --- a/include/clang/AST/OpenMPClause.h +++ b/include/clang/AST/OpenMPClause.h @@ -2629,7 +2629,6 @@ public: /// \param DepLoc Location of the dependency type. /// \param ColonLoc Colon location. /// \param VL List of references to the variables. - /// static OMPDependClause * Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc, OpenMPDependClauseKind DepKind, @@ -2648,6 +2647,14 @@ public: /// \brief Get colon location. SourceLocation getColonLoc() const { return ColonLoc; } + /// Set the loop counter value for the depend clauses with 'sink|source' kind + /// of dependency. Required for codegen. + void setCounterValue(Expr *V); + /// Get the loop counter value. + Expr *getCounterValue(); + /// Get the loop counter value. + const Expr *getCounterValue() const; + child_range children() { return child_range(reinterpret_cast(varlist_begin()), reinterpret_cast(varlist_end())); diff --git a/include/clang/AST/StmtOpenMP.h b/include/clang/AST/StmtOpenMP.h index 407bf76400..57e0294182 100644 --- a/include/clang/AST/StmtOpenMP.h +++ b/include/clang/AST/StmtOpenMP.h @@ -325,9 +325,10 @@ class OMPLoopDirective : public OMPExecutableDirective { EnsureUpperBoundOffset = 13, NextLowerBoundOffset = 14, NextUpperBoundOffset = 15, + NumIterationsOffset = 16, // Offset to the end (and start of the following counters/updates/finals // arrays) for worksharing loop directives. - WorksharingEnd = 16, + WorksharingEnd = 17, }; /// \brief Get the counters storage. @@ -475,6 +476,13 @@ protected: "expected worksharing loop directive"); *std::next(child_begin(), NextUpperBoundOffset) = NUB; } + void setNumIterations(Expr *NI) { + assert((isOpenMPWorksharingDirective(getDirectiveKind()) || + isOpenMPTaskLoopDirective(getDirectiveKind()) || + isOpenMPDistributeDirective(getDirectiveKind())) && + "expected worksharing loop directive"); + *std::next(child_begin(), NumIterationsOffset) = NI; + } void setCounters(ArrayRef A); void setPrivateCounters(ArrayRef A); void setInits(ArrayRef A); @@ -553,6 +561,7 @@ public: EUB = nullptr; NLB = nullptr; NUB = nullptr; + NumIterations = nullptr; Counters.resize(Size); PrivateCounters.resize(Size); Inits.resize(Size); @@ -660,6 +669,14 @@ public: return const_cast(reinterpret_cast( *std::next(child_begin(), NextUpperBoundOffset))); } + Expr *getNumIterations() const { + assert((isOpenMPWorksharingDirective(getDirectiveKind()) || + isOpenMPTaskLoopDirective(getDirectiveKind()) || + isOpenMPDistributeDirective(getDirectiveKind())) && + "expected worksharing loop directive"); + return const_cast(reinterpret_cast( + *std::next(child_begin(), NumIterationsOffset))); + } const Stmt *getBody() const { // This relies on the loop form is already checked by Sema. Stmt *Body = getAssociatedStmt()->IgnoreContainers(true); diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td index e9d12f0b18..c77244e40a 100644 --- a/include/clang/Basic/DiagnosticSemaKinds.td +++ b/include/clang/Basic/DiagnosticSemaKinds.td @@ -8239,8 +8239,6 @@ def err_omp_firstprivate_distribute_in_teams_reduction : Error< "reduction variable in '#pragma omp teams' cannot be firstprivate in '#pragma omp distribute'">; def err_omp_depend_clause_thread_simd : Error< "'depend' clauses cannot be mixed with '%0' clause">; -def err_omp_depend_sink_wrong_expr : Error< - "expected expression form x[+-d], where x is the loop iteration variable and d is a constant non-negative integer">; def err_omp_depend_sink_expected_loop_iteration : Error< "expected %0 loop iteration variable">; def err_omp_depend_sink_unexpected_expr : Error< diff --git a/lib/AST/OpenMPClause.cpp b/lib/AST/OpenMPClause.cpp index 73981cf0ff..01a2d11c54 100644 --- a/lib/AST/OpenMPClause.cpp +++ b/lib/AST/OpenMPClause.cpp @@ -498,7 +498,7 @@ OMPFlushClause *OMPFlushClause::Create(const ASTContext &C, SourceLocation LParenLoc, SourceLocation EndLoc, ArrayRef VL) { - void *Mem = C.Allocate(totalSizeToAlloc(VL.size())); + void *Mem = C.Allocate(totalSizeToAlloc(VL.size() + 1)); OMPFlushClause *Clause = new (Mem) OMPFlushClause(StartLoc, LParenLoc, EndLoc, VL.size()); Clause->setVarRefs(VL); @@ -510,26 +510,49 @@ OMPFlushClause *OMPFlushClause::CreateEmpty(const ASTContext &C, unsigned N) { return new (Mem) OMPFlushClause(N); } -OMPDependClause * -OMPDependClause::Create(const ASTContext &C, SourceLocation StartLoc, - SourceLocation LParenLoc, SourceLocation EndLoc, - OpenMPDependClauseKind DepKind, SourceLocation DepLoc, - SourceLocation ColonLoc, ArrayRef VL) { - void *Mem = C.Allocate(totalSizeToAlloc(VL.size())); +OMPDependClause *OMPDependClause::Create( + const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc, OpenMPDependClauseKind DepKind, + SourceLocation DepLoc, SourceLocation ColonLoc, ArrayRef VL) { + void *Mem = C.Allocate(totalSizeToAlloc(VL.size() + 1)); OMPDependClause *Clause = new (Mem) OMPDependClause(StartLoc, LParenLoc, EndLoc, VL.size()); Clause->setVarRefs(VL); Clause->setDependencyKind(DepKind); Clause->setDependencyLoc(DepLoc); Clause->setColonLoc(ColonLoc); + Clause->setCounterValue(nullptr); return Clause; } OMPDependClause *OMPDependClause::CreateEmpty(const ASTContext &C, unsigned N) { - void *Mem = C.Allocate(totalSizeToAlloc(N)); + void *Mem = C.Allocate(totalSizeToAlloc( + static_cast::type>(N) + + 1)); return new (Mem) OMPDependClause(N); } +void OMPDependClause::setCounterValue(Expr *V) { + assert(getDependencyKind() == OMPC_DEPEND_sink || + getDependencyKind() == OMPC_DEPEND_source || V == nullptr); + *getVarRefs().end() = V; +} + +const Expr *OMPDependClause::getCounterValue() const { + auto *V = *getVarRefs().end(); + assert(getDependencyKind() == OMPC_DEPEND_sink || + getDependencyKind() == OMPC_DEPEND_source || V == nullptr); + return V; +} + +Expr *OMPDependClause::getCounterValue() { + auto *V = *getVarRefs().end(); + assert(getDependencyKind() == OMPC_DEPEND_sink || + getDependencyKind() == OMPC_DEPEND_source || V == nullptr); + return V; +} + unsigned OMPClauseMappableExprCommon::getComponentsTotalNumber( MappableExprComponentListsRef ComponentLists) { unsigned TotalNum = 0u; diff --git a/lib/AST/StmtOpenMP.cpp b/lib/AST/StmtOpenMP.cpp index f3291a5113..b89e01c1cd 100644 --- a/lib/AST/StmtOpenMP.cpp +++ b/lib/AST/StmtOpenMP.cpp @@ -149,6 +149,7 @@ OMPForDirective::Create(const ASTContext &C, SourceLocation StartLoc, Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); @@ -199,6 +200,7 @@ OMPForSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc, Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); @@ -365,6 +367,7 @@ OMPParallelForDirective *OMPParallelForDirective::Create( Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); @@ -413,6 +416,7 @@ OMPParallelForSimdDirective *OMPParallelForSimdDirective::Create( Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); @@ -750,6 +754,7 @@ OMPTargetParallelForDirective *OMPTargetParallelForDirective::Create( Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); @@ -890,6 +895,7 @@ OMPTaskLoopDirective *OMPTaskLoopDirective::Create( Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); @@ -938,6 +944,7 @@ OMPTaskLoopSimdDirective *OMPTaskLoopSimdDirective::Create( Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); @@ -985,6 +992,7 @@ OMPDistributeDirective *OMPDistributeDirective::Create( Dir->setEnsureUpperBound(Exprs.EUB); Dir->setNextLowerBound(Exprs.NLB); Dir->setNextUpperBound(Exprs.NUB); + Dir->setNumIterations(Exprs.NumIterations); Dir->setCounters(Exprs.Counters); Dir->setPrivateCounters(Exprs.PrivateCounters); Dir->setInits(Exprs.Inits); diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp index 87e2ed61af..4b0d21373e 100644 --- a/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/lib/CodeGen/CGOpenMPRuntime.cpp @@ -627,6 +627,17 @@ enum OpenMPRTLFunction { // if_val, kmp_uint64 *lb, kmp_uint64 *ub, kmp_int64 st, int nogroup, int // sched, kmp_uint64 grainsize, void *task_dup); OMPRTL__kmpc_taskloop, + // Call to void __kmpc_doacross_init(ident_t *loc, kmp_int32 gtid, kmp_int32 + // num_dims, struct kmp_dim *dims); + OMPRTL__kmpc_doacross_init, + // Call to void __kmpc_doacross_fini(ident_t *loc, kmp_int32 gtid); + OMPRTL__kmpc_doacross_fini, + // Call to void __kmpc_doacross_post(ident_t *loc, kmp_int32 gtid, kmp_int64 + // *vec); + OMPRTL__kmpc_doacross_post, + // Call to void __kmpc_doacross_wait(ident_t *loc, kmp_int32 gtid, kmp_int64 + // *vec); + OMPRTL__kmpc_doacross_wait, // // Offloading related calls @@ -1476,6 +1487,46 @@ CGOpenMPRuntime::createRuntimeFunction(unsigned Function) { RTLFn = CGM.CreateRuntimeFunction(FnTy, /*Name=*/"__kmpc_taskloop"); break; } + case OMPRTL__kmpc_doacross_init: { + // Build void __kmpc_doacross_init(ident_t *loc, kmp_int32 gtid, kmp_int32 + // num_dims, struct kmp_dim *dims); + llvm::Type *TypeParams[] = {getIdentTyPointerTy(), + CGM.Int32Ty, + CGM.Int32Ty, + CGM.VoidPtrTy}; + llvm::FunctionType *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, /*Name=*/"__kmpc_doacross_init"); + break; + } + case OMPRTL__kmpc_doacross_fini: { + // Build void __kmpc_doacross_fini(ident_t *loc, kmp_int32 gtid); + llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty}; + llvm::FunctionType *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, /*Name=*/"__kmpc_doacross_fini"); + break; + } + case OMPRTL__kmpc_doacross_post: { + // Build void __kmpc_doacross_post(ident_t *loc, kmp_int32 gtid, kmp_int64 + // *vec); + llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty, + CGM.Int64Ty->getPointerTo()}; + llvm::FunctionType *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, /*Name=*/"__kmpc_doacross_post"); + break; + } + case OMPRTL__kmpc_doacross_wait: { + // Build void __kmpc_doacross_wait(ident_t *loc, kmp_int32 gtid, kmp_int64 + // *vec); + llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty, + CGM.Int64Ty->getPointerTo()}; + llvm::FunctionType *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, /*Name=*/"__kmpc_doacross_wait"); + break; + } case OMPRTL__tgt_target: { // Build int32_t __tgt_target(int32_t device_id, void *host_ptr, int32_t // arg_num, void** args_base, void **args, size_t *arg_sizes, int32_t @@ -6316,3 +6367,111 @@ void CGOpenMPRuntime::emitDeclareSimdFunction(const FunctionDecl *FD, emitX86DeclareSimdFunction(FD, Fn, VLENVal, ParamAttrs, State); } } + +namespace { +/// Cleanup action for doacross support. +class DoacrossCleanupTy final : public EHScopeStack::Cleanup { +public: + static const int DoacrossFinArgs = 2; + +private: + llvm::Value *RTLFn; + llvm::Value *Args[DoacrossFinArgs]; + +public: + DoacrossCleanupTy(llvm::Value *RTLFn, ArrayRef CallArgs) + : RTLFn(RTLFn) { + assert(CallArgs.size() == DoacrossFinArgs); + std::copy(CallArgs.begin(), CallArgs.end(), std::begin(Args)); + } + void Emit(CodeGenFunction &CGF, Flags /*flags*/) override { + if (!CGF.HaveInsertPoint()) + return; + CGF.EmitRuntimeCall(RTLFn, Args); + } +}; +} // namespace + +void CGOpenMPRuntime::emitDoacrossInit(CodeGenFunction &CGF, + const OMPLoopDirective &D) { + if (!CGF.HaveInsertPoint()) + return; + + ASTContext &C = CGM.getContext(); + QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/true); + RecordDecl *RD; + if (KmpDimTy.isNull()) { + // Build struct kmp_dim { // loop bounds info casted to kmp_int64 + // kmp_int64 lo; // lower + // kmp_int64 up; // upper + // kmp_int64 st; // stride + // }; + RD = C.buildImplicitRecord("kmp_dim"); + RD->startDefinition(); + addFieldToRecordDecl(C, RD, Int64Ty); + addFieldToRecordDecl(C, RD, Int64Ty); + addFieldToRecordDecl(C, RD, Int64Ty); + RD->completeDefinition(); + KmpDimTy = C.getRecordType(RD); + } else + RD = cast(KmpDimTy->getAsTagDecl()); + + Address DimsAddr = CGF.CreateMemTemp(KmpDimTy, "dims"); + CGF.EmitNullInitialization(DimsAddr, KmpDimTy); + enum { LowerFD = 0, UpperFD, StrideFD }; + // Fill dims with data. + LValue DimsLVal = CGF.MakeAddrLValue(DimsAddr, KmpDimTy); + // dims.upper = num_iterations; + LValue UpperLVal = + CGF.EmitLValueForField(DimsLVal, *std::next(RD->field_begin(), UpperFD)); + llvm::Value *NumIterVal = CGF.EmitScalarConversion( + CGF.EmitScalarExpr(D.getNumIterations()), D.getNumIterations()->getType(), + Int64Ty, D.getNumIterations()->getExprLoc()); + CGF.EmitStoreOfScalar(NumIterVal, UpperLVal); + // dims.stride = 1; + LValue StrideLVal = + CGF.EmitLValueForField(DimsLVal, *std::next(RD->field_begin(), StrideFD)); + CGF.EmitStoreOfScalar(llvm::ConstantInt::getSigned(CGM.Int64Ty, /*V=*/1), + StrideLVal); + + // Build call void __kmpc_doacross_init(ident_t *loc, kmp_int32 gtid, + // kmp_int32 num_dims, struct kmp_dim * dims); + llvm::Value *Args[] = {emitUpdateLocation(CGF, D.getLocStart()), + getThreadID(CGF, D.getLocStart()), + llvm::ConstantInt::getSigned(CGM.Int32Ty, 1), + CGF.Builder.CreatePointerBitCastOrAddrSpaceCast( + DimsAddr.getPointer(), CGM.VoidPtrTy)}; + + llvm::Value *RTLFn = createRuntimeFunction(OMPRTL__kmpc_doacross_init); + CGF.EmitRuntimeCall(RTLFn, Args); + llvm::Value *FiniArgs[DoacrossCleanupTy::DoacrossFinArgs] = { + emitUpdateLocation(CGF, D.getLocEnd()), getThreadID(CGF, D.getLocEnd())}; + llvm::Value *FiniRTLFn = createRuntimeFunction(OMPRTL__kmpc_doacross_fini); + CGF.EHStack.pushCleanup(NormalAndEHCleanup, FiniRTLFn, + llvm::makeArrayRef(FiniArgs)); +} + +void CGOpenMPRuntime::emitDoacrossOrdered(CodeGenFunction &CGF, + const OMPDependClause *C) { + QualType Int64Ty = + CGM.getContext().getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/1); + const Expr *CounterVal = C->getCounterValue(); + assert(CounterVal); + llvm::Value *CntVal = CGF.EmitScalarConversion(CGF.EmitScalarExpr(CounterVal), + CounterVal->getType(), Int64Ty, + CounterVal->getExprLoc()); + Address CntAddr = CGF.CreateMemTemp(Int64Ty, ".cnt.addr"); + CGF.EmitStoreOfScalar(CntVal, CntAddr, /*Volatile=*/false, Int64Ty); + llvm::Value *Args[] = {emitUpdateLocation(CGF, C->getLocStart()), + getThreadID(CGF, C->getLocStart()), + CntAddr.getPointer()}; + llvm::Value *RTLFn; + if (C->getDependencyKind() == OMPC_DEPEND_source) + RTLFn = createRuntimeFunction(OMPRTL__kmpc_doacross_post); + else { + assert(C->getDependencyKind() == OMPC_DEPEND_sink); + RTLFn = createRuntimeFunction(OMPRTL__kmpc_doacross_wait); + } + CGF.EmitRuntimeCall(RTLFn, Args); +} + diff --git a/lib/CodeGen/CGOpenMPRuntime.h b/lib/CodeGen/CGOpenMPRuntime.h index 63616f13f8..973e1f96b4 100644 --- a/lib/CodeGen/CGOpenMPRuntime.h +++ b/lib/CodeGen/CGOpenMPRuntime.h @@ -37,6 +37,7 @@ class Value; namespace clang { class Expr; class GlobalDecl; +class OMPDependClause; class OMPExecutableDirective; class OMPLoopDirective; class VarDecl; @@ -201,6 +202,12 @@ private: /// } flags; /// } kmp_depend_info_t; QualType KmpDependInfoTy; + /// struct kmp_dim { // loop bounds info casted to kmp_int64 + /// kmp_int64 lo; // lower + /// kmp_int64 up; // upper + /// kmp_int64 st; // stride + /// }; + QualType KmpDimTy; /// \brief Type struct __tgt_offload_entry{ /// void *addr; // Pointer to the offload entry info. /// // (function or global) @@ -1020,6 +1027,16 @@ public: /// attributes. virtual void emitDeclareSimdFunction(const FunctionDecl *FD, llvm::Function *Fn); + + /// Emit initialization for doacross loop nesting support. + /// \param D Loop-based construct used in doacross nesting construct. + virtual void emitDoacrossInit(CodeGenFunction &CGF, + const OMPLoopDirective &D); + + /// Emit code for doacross ordered directive with 'depend' clause. + /// \param C 'depend' clause with 'sink|source' dependency kind. + virtual void emitDoacrossOrdered(CodeGenFunction &CGF, + const OMPDependClause *C); }; } // namespace CodeGen diff --git a/lib/CodeGen/CGStmtOpenMP.cpp b/lib/CodeGen/CGStmtOpenMP.cpp index cfe4cb714d..45d8d0af10 100644 --- a/lib/CodeGen/CGStmtOpenMP.cpp +++ b/lib/CodeGen/CGStmtOpenMP.cpp @@ -1913,6 +1913,14 @@ bool CodeGenFunction::EmitOMPWorksharingLoop(const OMPLoopDirective &S) { incrementProfileCounter(&S); } + bool Ordered = false; + if (auto *OrderedClause = S.getSingleClause()) { + if (OrderedClause->getNumForLoops()) + RT.emitDoacrossInit(*this, S); + else + Ordered = true; + } + llvm::DenseSet EmittedFinals; emitAlignedClause(*this, S); EmitOMPLinearClauseInit(S); @@ -1960,7 +1968,6 @@ bool CodeGenFunction::EmitOMPWorksharingLoop(const OMPLoopDirective &S) { } const unsigned IVSize = getContext().getTypeSize(IVExpr->getType()); const bool IVSigned = IVExpr->getType()->hasSignedIntegerRepresentation(); - const bool Ordered = S.getSingleClause() != nullptr; // OpenMP 4.5, 2.7.1 Loop Construct, Description. // If the static schedule kind is specified or if the ordered clause is // specified, and if no monotonic modifier is specified, the effect will @@ -2685,8 +2692,11 @@ static llvm::Function *emitOutlinedOrderedFunction(CodeGenModule &CGM, } void CodeGenFunction::EmitOMPOrderedDirective(const OMPOrderedDirective &S) { - if (!S.getAssociatedStmt()) + if (!S.getAssociatedStmt()) { + for (const auto *DC : S.getClausesOfKind()) + CGM.getOpenMPRuntime().emitDoacrossOrdered(*this, DC); return; + } auto *C = S.getSingleClause(); auto &&CodeGen = [&S, C, this](CodeGenFunction &CGF, PrePostActionTy &Action) { diff --git a/lib/Sema/SemaExprMember.cpp b/lib/Sema/SemaExprMember.cpp index 0ce3974e51..c9231ed364 100644 --- a/lib/Sema/SemaExprMember.cpp +++ b/lib/Sema/SemaExprMember.cpp @@ -1792,6 +1792,7 @@ BuildFieldReferenceExpr(Sema &S, Expr *BaseExpr, bool IsArrow, // Build a reference to a private copy for non-static data members in // non-static member functions, privatized by OpenMP constructs. if (S.getLangOpts().OpenMP && IsArrow && + !S.CurContext->isDependentContext() && isa(Base.get()->IgnoreParenImpCasts())) { if (auto *PrivateCopy = S.IsOpenMPCapturedDecl(Field)) return S.getOpenMPCapturedExpr(PrivateCopy, VK, OK, OpLoc); diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp index f249e3f1e0..57759bcb63 100644 --- a/lib/Sema/SemaOpenMP.cpp +++ b/lib/Sema/SemaOpenMP.cpp @@ -57,6 +57,8 @@ public: SourceLocation ImplicitDSALoc; DSAVarData() {} }; + typedef llvm::SmallVector, 4> + OperatorOffsetTy; private: struct DSAInfo final { @@ -75,6 +77,8 @@ private: MappedExprComponentsTy; typedef llvm::StringMap> CriticalsWithHintsTy; + typedef llvm::DenseMap + DoacrossDependMapTy; struct SharingMapTy final { DeclSAMapTy SharingMap; @@ -87,6 +91,10 @@ private: DeclarationNameInfo DirectiveName; Scope *CurScope = nullptr; SourceLocation ConstructLoc; + /// Set of 'depend' clauses with 'sink|source' dependence kind. Required to + /// get the data (loop counters etc.) about enclosing loop-based construct. + /// This data is required during codegen. + DoacrossDependMapTy DoacrossDepends; /// \brief first argument (Expr *) contains optional argument of the /// 'ordered' clause, the second one is true if the regions has 'ordered' /// clause, false otherwise. @@ -360,6 +368,21 @@ public: assert(Stack.size() > 1); return Stack.size() - 2; } + void addDoacrossDependClause(OMPDependClause *C, OperatorOffsetTy &OpsOffs) { + assert(Stack.size() > 2); + assert(isOpenMPWorksharingDirective(Stack[Stack.size() - 2].Directive)); + Stack[Stack.size() - 2].DoacrossDepends.insert({C, OpsOffs}); + } + llvm::iterator_range + getDoacrossDependClauses() const { + assert(Stack.size() > 1); + if (isOpenMPWorksharingDirective(Stack[Stack.size() - 1].Directive)) { + auto &Ref = Stack[Stack.size() - 1].DoacrossDepends; + return llvm::make_range(Ref.begin(), Ref.end()); + } + return llvm::make_range(Stack[0].DoacrossDepends.end(), + Stack[0].DoacrossDepends.end()); + } }; bool isParallelOrTaskRegion(OpenMPDirectiveKind DKind) { return isOpenMPParallelDirective(DKind) || isOpenMPTaskingDirective(DKind) || @@ -4190,23 +4213,23 @@ Expr *OpenMPIterationSpaceChecker::BuildCounterInit() const { return LB; } Expr *OpenMPIterationSpaceChecker::BuildCounterStep() const { return Step; } /// \brief Iteration space of a single for loop. -struct LoopIterationSpace { +struct LoopIterationSpace final { /// \brief Condition of the loop. - Expr *PreCond; + Expr *PreCond = nullptr; /// \brief This expression calculates the number of iterations in the loop. /// It is always possible to calculate it before starting the loop. - Expr *NumIterations; + Expr *NumIterations = nullptr; /// \brief The loop counter variable. - Expr *CounterVar; + Expr *CounterVar = nullptr; /// \brief Private loop counter variable. - Expr *PrivateCounterVar; + Expr *PrivateCounterVar = nullptr; /// \brief This is initializer for the initial value of #CounterVar. - Expr *CounterInit; + Expr *CounterInit = nullptr; /// \brief This is step for the #CounterVar used to generate its update: /// #CounterVar = #CounterInit + #CounterStep * CurrentIteration. - Expr *CounterStep; + Expr *CounterStep = nullptr; /// \brief Should step be subtracted? - bool Subtract; + bool Subtract = false; /// \brief Source range of the loop init. SourceRange InitSrcRange; /// \brief Source range of the loop condition. @@ -4864,6 +4887,7 @@ CheckOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr, Built.Inits.resize(NestedLoopCount); Built.Updates.resize(NestedLoopCount); Built.Finals.resize(NestedLoopCount); + SmallVector LoopMultipliers; { ExprResult Div; // Go from inner nested loop to outer. @@ -4928,11 +4952,12 @@ CheckOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr, // Add parentheses (for debugging purposes only). if (Div.isUsable()) - Div = SemaRef.ActOnParenExpr(UpdLoc, UpdLoc, Div.get()); + Div = tryBuildCapture(SemaRef, Div.get(), Captures); if (!Div.isUsable()) { HasErrors = true; break; } + LoopMultipliers.push_back(Div.get()); } if (!Update.isUsable() || !Final.isUsable()) { HasErrors = true; @@ -4969,6 +4994,54 @@ CheckOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr, Built.NLB = NextLB.get(); Built.NUB = NextUB.get(); + Expr *CounterVal = SemaRef.DefaultLvalueConversion(IV.get()).get(); + // Fill data for doacross depend clauses. + for (auto Pair : DSA.getDoacrossDependClauses()) { + if (Pair.first->getDependencyKind() == OMPC_DEPEND_source) + Pair.first->setCounterValue(CounterVal); + else { + if (NestedLoopCount != Pair.second.size() || + NestedLoopCount != LoopMultipliers.size() + 1) { + // Erroneous case - clause has some problems. + Pair.first->setCounterValue(CounterVal); + continue; + } + assert(Pair.first->getDependencyKind() == OMPC_DEPEND_sink); + auto I = Pair.second.rbegin(); + auto IS = IterSpaces.rbegin(); + auto ILM = LoopMultipliers.rbegin(); + Expr *UpCounterVal = CounterVal; + Expr *Multiplier = nullptr; + for (int Cnt = NestedLoopCount - 1; Cnt >= 0; --Cnt) { + if (I->first) { + assert(IS->CounterStep); + Expr *NormalizedOffset = + SemaRef + .BuildBinOp(CurScope, I->first->getExprLoc(), BO_Div, + I->first, IS->CounterStep) + .get(); + if (Multiplier) { + NormalizedOffset = + SemaRef + .BuildBinOp(CurScope, I->first->getExprLoc(), BO_Mul, + NormalizedOffset, Multiplier) + .get(); + } + assert(I->second == OO_Plus || I->second == OO_Minus); + BinaryOperatorKind BOK = (I->second == OO_Plus) ? BO_Add : BO_Sub; + UpCounterVal = + SemaRef.BuildBinOp(CurScope, I->first->getExprLoc(), BOK, + UpCounterVal, NormalizedOffset).get(); + } + Multiplier = *ILM; + ++I; + ++IS; + ++ILM; + } + Pair.first->setCounterValue(UpCounterVal); + } + } + return NestedLoopCount; } @@ -9484,6 +9557,7 @@ Sema::ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, return nullptr; } SmallVector Vars; + DSAStackTy::OperatorOffsetTy OpsOffs; llvm::APSInt DepCounter(/*BitWidth=*/32); llvm::APSInt TotalDepCount(/*BitWidth=*/32); if (DepKind == OMPC_DEPEND_sink) { @@ -9496,8 +9570,7 @@ Sema::ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, DSAStack->getParentOrderedRegionParam()) { for (auto &RefExpr : VarList) { assert(RefExpr && "NULL expr in OpenMP shared clause."); - if (isa(RefExpr) || - (DepKind == OMPC_DEPEND_sink && CurContext->isDependentContext())) { + if (isa(RefExpr)) { // It will be analyzed later. Vars.push_back(RefExpr); continue; @@ -9519,62 +9592,66 @@ Sema::ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, // directive, xi denotes the loop iteration variable of the i-th nested // loop associated with the loop directive, and di is a constant // non-negative integer. + if (CurContext->isDependentContext()) { + // It will be analyzed later. + Vars.push_back(RefExpr); + continue; + } SimpleExpr = SimpleExpr->IgnoreImplicit(); - auto *DE = dyn_cast(SimpleExpr); - if (!DE) { - OverloadedOperatorKind OOK = OO_None; - SourceLocation OOLoc; - Expr *LHS, *RHS; - if (auto *BO = dyn_cast(SimpleExpr)) { - OOK = BinaryOperator::getOverloadedOperator(BO->getOpcode()); - OOLoc = BO->getOperatorLoc(); - LHS = BO->getLHS()->IgnoreParenImpCasts(); - RHS = BO->getRHS()->IgnoreParenImpCasts(); - } else if (auto *OCE = dyn_cast(SimpleExpr)) { - OOK = OCE->getOperator(); - OOLoc = OCE->getOperatorLoc(); - LHS = OCE->getArg(/*Arg=*/0)->IgnoreParenImpCasts(); - RHS = OCE->getArg(/*Arg=*/1)->IgnoreParenImpCasts(); - } else if (auto *MCE = dyn_cast(SimpleExpr)) { - OOK = MCE->getMethodDecl() - ->getNameInfo() - .getName() - .getCXXOverloadedOperator(); - OOLoc = MCE->getCallee()->getExprLoc(); - LHS = MCE->getImplicitObjectArgument()->IgnoreParenImpCasts(); - RHS = MCE->getArg(/*Arg=*/0)->IgnoreParenImpCasts(); - } else { - Diag(ELoc, diag::err_omp_depend_sink_wrong_expr); - continue; - } - DE = dyn_cast(LHS); - if (!DE) { - Diag(LHS->getExprLoc(), - diag::err_omp_depend_sink_expected_loop_iteration) - << DSAStack->getParentLoopControlVariable( - DepCounter.getZExtValue()); - continue; - } - if (OOK != OO_Plus && OOK != OO_Minus) { - Diag(OOLoc, diag::err_omp_depend_sink_expected_plus_minus); - continue; - } - ExprResult Res = VerifyPositiveIntegerConstantInClause( + OverloadedOperatorKind OOK = OO_None; + SourceLocation OOLoc; + Expr *LHS = SimpleExpr; + Expr *RHS = nullptr; + if (auto *BO = dyn_cast(SimpleExpr)) { + OOK = BinaryOperator::getOverloadedOperator(BO->getOpcode()); + OOLoc = BO->getOperatorLoc(); + LHS = BO->getLHS()->IgnoreParenImpCasts(); + RHS = BO->getRHS()->IgnoreParenImpCasts(); + } else if (auto *OCE = dyn_cast(SimpleExpr)) { + OOK = OCE->getOperator(); + OOLoc = OCE->getOperatorLoc(); + LHS = OCE->getArg(/*Arg=*/0)->IgnoreParenImpCasts(); + RHS = OCE->getArg(/*Arg=*/1)->IgnoreParenImpCasts(); + } else if (auto *MCE = dyn_cast(SimpleExpr)) { + OOK = MCE->getMethodDecl() + ->getNameInfo() + .getName() + .getCXXOverloadedOperator(); + OOLoc = MCE->getCallee()->getExprLoc(); + LHS = MCE->getImplicitObjectArgument()->IgnoreParenImpCasts(); + RHS = MCE->getArg(/*Arg=*/0)->IgnoreParenImpCasts(); + } + SourceLocation ELoc; + SourceRange ERange; + auto Res = getPrivateItem(*this, LHS, ELoc, ERange, + /*AllowArraySection=*/false); + if (Res.second) { + // It will be analyzed later. + Vars.push_back(RefExpr); + } + ValueDecl *D = Res.first; + if (!D) + continue; + + if (OOK != OO_Plus && OOK != OO_Minus && (RHS || OOK != OO_None)) { + Diag(OOLoc, diag::err_omp_depend_sink_expected_plus_minus); + continue; + } + if (RHS) { + ExprResult RHSRes = VerifyPositiveIntegerConstantInClause( RHS, OMPC_depend, /*StrictlyPositive=*/false); - if (Res.isInvalid()) + if (RHSRes.isInvalid()) continue; } - auto *VD = dyn_cast(DE->getDecl()); if (!CurContext->isDependentContext() && DSAStack->getParentOrderedRegionParam() && - (!VD || - DepCounter != DSAStack->isParentLoopControlVariable(VD).first)) { - Diag(DE->getExprLoc(), - diag::err_omp_depend_sink_expected_loop_iteration) + DepCounter != DSAStack->isParentLoopControlVariable(D).first) { + Diag(ELoc, diag::err_omp_depend_sink_expected_loop_iteration) << DSAStack->getParentLoopControlVariable( - DepCounter.getZExtValue()); + DepCounter.getZExtValue()); continue; } + OpsOffs.push_back({RHS, OOK}); } else { // OpenMP [2.11.1.1, Restrictions, p.3] // A variable that is part of another variable (such as a field of a @@ -9596,7 +9673,6 @@ Sema::ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, continue; } } - Vars.push_back(RefExpr->IgnoreParenImpCasts()); } @@ -9610,9 +9686,11 @@ Sema::ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, Vars.empty()) return nullptr; } - - return OMPDependClause::Create(Context, StartLoc, LParenLoc, EndLoc, DepKind, - DepLoc, ColonLoc, Vars); + auto *C = OMPDependClause::Create(Context, StartLoc, LParenLoc, EndLoc, + DepKind, DepLoc, ColonLoc, Vars); + if (DepKind == OMPC_DEPEND_sink || DepKind == OMPC_DEPEND_source) + DSAStack->addDoacrossDependClause(C, OpsOffs); + return C; } OMPClause *Sema::ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc, diff --git a/lib/Serialization/ASTReaderStmt.cpp b/lib/Serialization/ASTReaderStmt.cpp index 3d6b8aceda..1d64753259 100644 --- a/lib/Serialization/ASTReaderStmt.cpp +++ b/lib/Serialization/ASTReaderStmt.cpp @@ -2215,6 +2215,7 @@ void OMPClauseReader::VisitOMPDependClause(OMPDependClause *C) { for (unsigned i = 0; i != NumVars; ++i) Vars.push_back(Reader->Reader.ReadSubExpr()); C->setVarRefs(Vars); + C->setCounterValue(Reader->Reader.ReadSubExpr()); } void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) { @@ -2360,6 +2361,7 @@ void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) { D->setEnsureUpperBound(Reader.ReadSubExpr()); D->setNextLowerBound(Reader.ReadSubExpr()); D->setNextUpperBound(Reader.ReadSubExpr()); + D->setNumIterations(Reader.ReadSubExpr()); } SmallVector Sub; unsigned CollapsedNum = D->getCollapsedNumber(); diff --git a/lib/Serialization/ASTWriterStmt.cpp b/lib/Serialization/ASTWriterStmt.cpp index 929faccbad..35d8b92ded 100644 --- a/lib/Serialization/ASTWriterStmt.cpp +++ b/lib/Serialization/ASTWriterStmt.cpp @@ -2012,6 +2012,7 @@ void OMPClauseWriter::VisitOMPDependClause(OMPDependClause *C) { Record.AddSourceLocation(C->getColonLoc()); for (auto *VE : C->varlists()) Record.AddStmt(VE); + Record.AddStmt(C->getCounterValue()); } void OMPClauseWriter::VisitOMPDeviceClause(OMPDeviceClause *C) { @@ -2127,6 +2128,7 @@ void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) { Record.AddStmt(D->getEnsureUpperBound()); Record.AddStmt(D->getNextLowerBound()); Record.AddStmt(D->getNextUpperBound()); + Record.AddStmt(D->getNumIterations()); } for (auto I : D->counters()) { Record.AddStmt(I); diff --git a/test/OpenMP/ordered_doacross_codegen.cpp b/test/OpenMP/ordered_doacross_codegen.cpp new file mode 100644 index 0000000000..d1fe99d4b8 --- /dev/null +++ b/test/OpenMP/ordered_doacross_codegen.cpp @@ -0,0 +1,124 @@ +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s +// REQUIRES: x86-registered-target +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +// CHECK: [[KMP_DIM:%.+]] = type { i64, i64, i64 } +extern int n; +int a[10], b[10], c[10], d[10]; +void foo(); + +// CHECK-LABEL: @main() +int main() { + int i; +// CHECK: [[DIMS:%.+]] = alloca [[KMP_DIM]], +// CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num([[IDENT:%.+]]) +// CHECK: icmp +// CHECK-NEXT: br i1 % +// CHECK: [[CAST:%.+]] = bitcast [[KMP_DIM]]* [[DIMS]] to i8* +// CHECK: call void @llvm.memset.p0i8.i64(i8* [[CAST]], i8 0, i64 24, i32 8, i1 false) +// CHECK: getelementptr inbounds [[KMP_DIM]], [[KMP_DIM]]* [[DIMS]], i32 0, i32 1 +// CHECK: store i64 %{{.+}}, i64* % +// CHECK: getelementptr inbounds [[KMP_DIM]], [[KMP_DIM]]* [[DIMS]], i32 0, i32 2 +// CHECK: store i64 1, i64* % +// CHECK: [[CAST:%.+]] = bitcast [[KMP_DIM]]* [[DIMS]] to i8* +// CHECK: call void @__kmpc_doacross_init([[IDENT]], i32 [[GTID]], i32 1, i8* [[CAST]]) +// CHECK: call void @__kmpc_for_static_init_4( +#pragma omp for ordered(1) + for (i = 0; i < n; ++i) { + a[i] = b[i] + 1; + foo(); +// CHECK: invoke void [[FOO:.+]]( +// CHECK: load i32, i32* [[CNT:%.+]], +// CHECK-NEXT: sext i32 %{{.+}} to i64 +// CHECK-NEXT: store i64 %{{.+}}, i64* [[TMP:%.+]], +// CHECK-NEXT: call void @__kmpc_doacross_post([[IDENT]], i32 [[GTID]], i64* [[TMP]]) +#pragma omp ordered depend(source) + c[i] = c[i] + 1; + foo(); +// CHECK: invoke void [[FOO]] +// CHECK: load i32, i32* [[CNT]], +// CHECK-NEXT: sub nsw i32 %{{.+}}, 2 +// CHECK-NEXT: sext i32 %{{.+}} to i64 +// CHECK-NEXT: store i64 %{{.+}}, i64* [[TMP:%.+]], +// CHECK-NEXT: call void @__kmpc_doacross_wait([[IDENT]], i32 [[GTID]], i64* [[TMP]]) +#pragma omp ordered depend(sink : i - 2) + d[i] = a[i - 2]; + } + // CHECK: landingpad + // CHECK: call void @__kmpc_doacross_fini([[IDENT]], i32 [[GTID]]) + // CHECK: br label % + + // CHECK: call void @__kmpc_for_static_fini( + // CHECK: call void @__kmpc_doacross_fini([[IDENT]], i32 [[GTID]]) + // CHECK: ret i32 0 + return 0; +} + +// CHECK: define {{.+}}TestStruct +template +struct TestStruct { + static const int M = 10; + static const int N = 20; + T i; + T a[N][M]; + T b[N][M]; + T foo(T, T); + T bar(T, T, T); + void baz(T, T); + TestStruct() { +// CHECK: [[CNT:%.+]] = alloca i64, +// CHECK: [[DIMS:%.+]] = alloca [[KMP_DIM]], +// CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num([[IDENT:%.+]]) +// CHECK: icmp +// CHECK-NEXT: br i1 % +// CHECK: [[CAST:%.+]] = bitcast [[KMP_DIM]]* [[DIMS]] to i8* +// CHECK: call void @llvm.memset.p0i8.i64(i8* [[CAST]], i8 0, i64 24, i32 8, i1 false) +// CHECK: getelementptr inbounds [[KMP_DIM]], [[KMP_DIM]]* [[DIMS]], i32 0, i32 1 +// CHECK: store i64 %{{.+}}, i64* % +// CHECK: getelementptr inbounds [[KMP_DIM]], [[KMP_DIM]]* [[DIMS]], i32 0, i32 2 +// CHECK: store i64 1, i64* % +// CHECK: [[CAST:%.+]] = bitcast [[KMP_DIM]]* [[DIMS]] to i8* +// CHECK: call void @__kmpc_doacross_init([[IDENT]], i32 [[GTID]], i32 1, i8* [[CAST]]) +// CHECK: call void @__kmpc_for_static_init_8( +#pragma omp for ordered(2) + for (T j = 0; j < M; j++) + for (i = 0; i < n; i += 2) { + a[i][j] = foo(i, j); +// CHECK: invoke {{.+TestStruct.+foo}} +// CHECK: load i64, i64* [[CNT]], +// CHECK-NEXT: sub nsw i64 %{{.+}}, 1 +// CHECK-NEXT: store i64 %{{.+}}, i64* [[TMP:%.+]], +// CHECK-NEXT: call void @__kmpc_doacross_wait([[IDENT]], i32 [[GTID]], i64* [[TMP]]) +// CHECK-NEXT: load i64, i64* [[CNT]], +// CHECK-NEXT: load i32, i32* % +// CHECK-NEXT: mul nsw i32 1, % +// CHECK-NEXT: sext i32 %{{.+}} to i64 +// CHECK-NEXT: sub nsw i64 % +// CHECK-NEXT: store i64 %{{.+}}, i64* [[TMP:%.+]], +// CHECK-NEXT: call void @__kmpc_doacross_wait([[IDENT]], i32 [[GTID]], i64* [[TMP]]) +#pragma omp ordered depend(sink : j, i - 2) depend(sink : j - 1, i) + b[i][j] = bar(a[i][j], b[i - 1][j], b[i][j - 1]); +// CHECK: invoke {{.+TestStruct.+bar}} +// CHECK: load i64, i64* [[CNT]], +// CHECK-NEXT: store i64 %{{.+}}, i64* [[TMP:%.+]], +// CHECK-NEXT: call void @__kmpc_doacross_post([[IDENT]], i32 [[GTID]], i64* [[TMP]]) +#pragma omp ordered depend(source) + baz(a[i][j], b[i][j]); + } + } + // CHECK: landingpad + // CHECK: call void @__kmpc_doacross_fini([[IDENT]], i32 [[GTID]]) + // CHECK: br label % + + // CHECK: call void @__kmpc_for_static_fini( + // CHECK: call void @__kmpc_doacross_fini([[IDENT]], i32 [[GTID]]) + // CHECK: ret +}; + +TestStruct s; +#endif // HEADER