From ef9ec4bbcca2fa4f64df47bc426f1d1c59ea47e2 Mon Sep 17 00:00:00 2001 From: Johannes Doerfert Date: Thu, 2 Mar 2023 18:35:15 -0800 Subject: [PATCH] [OpenMP] Add the `ompx_attribute` clause for target directives CUDA and HIP have kernel attributes to tune the code generation (in the backend). To reuse this functionality for OpenMP target regions we introduce the `ompx_attribute` clause that takes these kernel attributes and emits code as if they had been attached to the kernel fuction (which is implicitly generated). To limit the impact, we only support three kernel attributes: `amdgpu_waves_per_eu`, for AMDGPU `amdgpu_flat_work_group_size`, for AMDGPU `launch_bounds`, for NVPTX The existing implementations of those attributes are used for error checking and code generation. `ompx_attribute` can be attached to any executable target region and it can hold more than one kernel attribute. Differential Revision: https://reviews.llvm.org/D156184 --- clang/include/clang/AST/OpenMPClause.h | 48 ++++++++++ clang/include/clang/AST/RecursiveASTVisitor.h | 6 ++ clang/include/clang/Basic/DiagnosticGroups.td | 3 +- .../clang/Basic/DiagnosticParseKinds.td | 3 + clang/include/clang/Parse/Parser.h | 7 ++ clang/include/clang/Sema/Sema.h | 21 +++++ clang/lib/AST/OpenMPClause.cpp | 12 +++ clang/lib/AST/StmtProfile.cpp | 2 + clang/lib/CodeGen/CGOpenMPRuntime.cpp | 19 +++- clang/lib/CodeGen/CodeGenModule.h | 15 +++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 82 +++++++++-------- clang/lib/CodeGen/Targets/NVPTX.cpp | 44 +++++---- clang/lib/Parse/ParseOpenMP.cpp | 59 ++++++++++++ clang/lib/Sema/SemaDeclAttr.cpp | 53 +++++++---- clang/lib/Sema/SemaOpenMP.cpp | 8 ++ clang/lib/Sema/TreeTransform.h | 22 +++++ clang/lib/Serialization/ASTReader.cpp | 12 +++ clang/lib/Serialization/ASTWriter.cpp | 7 ++ clang/test/OpenMP/ompx_attributes_codegen.cpp | 31 +++++++ .../test/OpenMP/ompx_attributes_messages.cpp | 47 ++++++++++ clang/tools/libclang/CIndex.cpp | 2 + llvm/include/llvm/Frontend/OpenMP/OMP.td | 91 +++++++++++++------ 22 files changed, 489 insertions(+), 105 deletions(-) create mode 100644 clang/test/OpenMP/ompx_attributes_codegen.cpp create mode 100644 clang/test/OpenMP/ompx_attributes_messages.cpp diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h index 0bea21270692..31ae3d42e232 100644 --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -9172,6 +9172,54 @@ public: } }; +/// This represents 'ompx_attribute' clause in a directive that might generate +/// an outlined function. An example is given below. +/// +/// \code +/// #pragma omp target [...] ompx_attribute(flatten) +/// \endcode +class OMPXAttributeClause + : public OMPNoChildClause { + friend class OMPClauseReader; + + /// Location of '('. + SourceLocation LParenLoc; + + /// The parsed attributes (clause arguments) + SmallVector Attrs; + +public: + /// Build 'ompx_attribute' clause. + /// + /// \param Attrs The parsed attributes (clause arguments) + /// \param StartLoc Starting location of the clause. + /// \param LParenLoc Location of '('. + /// \param EndLoc Ending location of the clause. + OMPXAttributeClause(ArrayRef Attrs, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc) + : OMPNoChildClause(StartLoc, EndLoc), LParenLoc(LParenLoc), Attrs(Attrs) { + } + + /// Build an empty clause. + OMPXAttributeClause() : OMPNoChildClause() {} + + /// Sets the location of '('. + void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; } + + /// Returns the location of '('. + SourceLocation getLParenLoc() const { return LParenLoc; } + + /// Returned the attributes parsed from this clause. + ArrayRef getAttrs() const { return Attrs; } + +private: + /// Replace the attributes with \p NewAttrs. + void setAttrs(ArrayRef NewAttrs) { + Attrs.clear(); + Attrs.append(NewAttrs.begin(), NewAttrs.end()); + } +}; + } // namespace clang #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h index 604875cd6337..fc2d1ff708bf 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -3875,6 +3875,12 @@ bool RecursiveASTVisitor::VisitOMPDoacrossClause( return true; } +template +bool RecursiveASTVisitor::VisitOMPXAttributeClause( + OMPXAttributeClause *C) { + return true; +} + // FIXME: look at the following tricky-seeming exprs to see if we // need to recurse on anything. These are ones that have methods // returning decls or qualtypes or nestednamespecifier -- though I'm diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td index 7b4d415bf064..6a0a01e4a981 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -1278,9 +1278,10 @@ def OpenMPMapping : DiagGroup<"openmp-mapping">; def OpenMPTarget : DiagGroup<"openmp-target", [OpenMPMapping]>; def OpenMPPre51Compat : DiagGroup<"pre-openmp-51-compat">; def OpenMP51Ext : DiagGroup<"openmp-51-extensions">; +def OpenMPExtensions : DiagGroup<"openmp-extensions">; def OpenMP : DiagGroup<"openmp", [ SourceUsesOpenMP, OpenMPClauses, OpenMPLoopForm, OpenMPTarget, - OpenMPMapping, OpenMP51Ext + OpenMPMapping, OpenMP51Ext, OpenMPExtensions ]>; // Backend warnings. diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td index 8d729c31641e..a804442ab34e 100644 --- a/clang/include/clang/Basic/DiagnosticParseKinds.td +++ b/clang/include/clang/Basic/DiagnosticParseKinds.td @@ -1540,6 +1540,9 @@ def warn_omp_more_one_omp_all_memory : Warning< InGroup; def warn_omp_depend_in_ordered_deprecated : Warning<"'depend' clause for" " 'ordered' is deprecated; use 'doacross' instead">, InGroup; +def warn_omp_invalid_attribute_for_ompx_attributes : Warning<"'ompx_attribute' clause only allows " + "'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; " + "%0 is ignored">, InGroup; // Pragma loop support. def err_pragma_loop_missing_argument : Error< diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h index 475dfe845528..b58041477c9e 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3490,6 +3490,13 @@ private: // OMPClause *ParseOpenMPInteropClause(OpenMPClauseKind Kind, bool ParseOnly); + /// Parses a ompx_attribute clause + /// + /// \param ParseOnly true to skip the clause's semantic actions and return + /// nullptr. + // + OMPClause *ParseOpenMPOMPXAttributesClause(bool ParseOnly); + public: /// Parses simple expression in parens for single-expression clauses of OpenMP /// constructs. diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 3418a37b3077..7c641d5e273a 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10988,6 +10988,11 @@ public: bool ConstantFoldAttrArgs(const AttributeCommonInfo &CI, MutableArrayRef Args); + /// Create an CUDALaunchBoundsAttr attribute. + CUDALaunchBoundsAttr *CreateLaunchBoundsAttr(const AttributeCommonInfo &CI, + Expr *MaxThreads, + Expr *MinBlocks); + /// AddLaunchBoundsAttr - Adds a launch_bounds attribute to a particular /// declaration. void AddLaunchBoundsAttr(Decl *D, const AttributeCommonInfo &CI, @@ -11004,11 +11009,21 @@ public: void AddXConsumedAttr(Decl *D, const AttributeCommonInfo &CI, RetainOwnershipKind K, bool IsTemplateInstantiation); + /// Create an AMDGPUWavesPerEUAttr attribute. + AMDGPUFlatWorkGroupSizeAttr * + CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min, + Expr *Max); + /// addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size /// attribute to a particular declaration. void addAMDGPUFlatWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max); + /// Create an AMDGPUWavesPerEUAttr attribute. + AMDGPUWavesPerEUAttr * + CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI, Expr *Min, + Expr *Max); + /// addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a /// particular declaration. void addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, @@ -12341,6 +12356,12 @@ public: ArrayRef VarList, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); + /// Called on a well-formed 'ompx_attribute' clause. + OMPClause *ActOnOpenMPXAttributeClause(ArrayRef Attrs, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc); + /// The kind of conversion being performed. enum CheckedConversionKind { /// An implicit conversion. diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp index 4c895822ffdf..f5ad75028a64 100644 --- a/clang/lib/AST/OpenMPClause.cpp +++ b/clang/lib/AST/OpenMPClause.cpp @@ -2534,6 +2534,18 @@ void OMPClausePrinter::VisitOMPDoacrossClause(OMPDoacrossClause *Node) { OS << ")"; } +void OMPClausePrinter::VisitOMPXAttributeClause(OMPXAttributeClause *Node) { + OS << "ompx_attribute("; + bool IsFirst = true; + for (auto &Attr : Node->getAttrs()) { + if (!IsFirst) + OS << ", "; + Attr->printPretty(OS, Policy); + IsFirst = false; + } + OS << ")"; +} + void OMPTraitInfo::getAsVariantMatchInfo(ASTContext &ASTCtx, VariantMatchInfo &VMI) const { for (const OMPTraitSet &Set : Sets) { diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index d8a667b2d0fd..60646f7a0da5 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -928,6 +928,8 @@ void OMPClauseProfiler::VisitOMPXDynCGroupMemClause( void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) { VisitOMPClauseList(C); } +void OMPClauseProfiler::VisitOMPXAttributeClause(const OMPXAttributeClause *C) { +} } // namespace void diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index a52ec8909b12..a4f7eb96c0d9 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6110,8 +6110,23 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper( DefaultValTeams, DefaultValThreads, IsOffloadEntry, OutlinedFn, OutlinedFnID); - if (OutlinedFn != nullptr) - CGM.getTargetCodeGenInfo().setTargetAttributes(nullptr, OutlinedFn, CGM); + if (!OutlinedFn) + return; + + CGM.getTargetCodeGenInfo().setTargetAttributes(nullptr, OutlinedFn, CGM); + + for (auto *C : D.getClausesOfKind()) { + for (auto *A : C->getAttrs()) { + if (auto *Attr = dyn_cast(A)) + CGM.handleCUDALaunchBoundsAttr(OutlinedFn, Attr); + else if (auto *Attr = dyn_cast(A)) + CGM.handleAMDGPUFlatWorkGroupSizeAttr(OutlinedFn, Attr); + else if (auto *Attr = dyn_cast(A)) + CGM.handleAMDGPUWavesPerEUAttr(OutlinedFn, Attr); + else + llvm_unreachable("Unexpected attribute kind"); + } + } } /// Checks if the expression is constant or does not have non-trivial function diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index 05cb217e2bee..f5fd94492540 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1557,6 +1557,21 @@ public: /// because we'll lose all important information after each repl. void moveLazyEmissionStates(CodeGenModule *NewBuilder); + /// Emit the IR encoding to attach the CUDA launch bounds attribute to \p F. + void handleCUDALaunchBoundsAttr(llvm::Function *F, + const CUDALaunchBoundsAttr *A); + + /// Emit the IR encoding to attach the AMD GPU flat-work-group-size attribute + /// to \p F. Alternatively, the work group size can be taken from a \p + /// ReqdWGS. + void handleAMDGPUFlatWorkGroupSizeAttr( + llvm::Function *F, const AMDGPUFlatWorkGroupSizeAttr *A, + const ReqdWorkGroupSizeAttr *ReqdWGS = nullptr); + + /// Emit the IR encoding to attach the AMD GPU waves-per-eu attribute to \p F. + void handleAMDGPUWavesPerEUAttr(llvm::Function *F, + const AMDGPUWavesPerEUAttr *A); + private: llvm::Constant *GetOrCreateLLVMFunction( StringRef MangledName, llvm::Type *Ty, GlobalDecl D, bool ForVTable, diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 796a2be81a09..bac7787643e3 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -317,26 +317,7 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( const auto *FlatWGS = FD->getAttr(); if (ReqdWGS || FlatWGS) { - unsigned Min = 0; - unsigned Max = 0; - if (FlatWGS) { - Min = FlatWGS->getMin() - ->EvaluateKnownConstInt(M.getContext()) - .getExtValue(); - Max = FlatWGS->getMax() - ->EvaluateKnownConstInt(M.getContext()) - .getExtValue(); - } - if (ReqdWGS && Min == 0 && Max == 0) - Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim(); - - if (Min != 0) { - assert(Min <= Max && "Min must be less than or equal Max"); - - std::string AttrVal = llvm::utostr(Min) + "," + llvm::utostr(Max); - F->addFnAttr("amdgpu-flat-work-group-size", AttrVal); - } else - assert(Max == 0 && "Max must be zero"); + M.handleAMDGPUFlatWorkGroupSizeAttr(F, FlatWGS, ReqdWGS); } else if (IsOpenCLKernel || IsHIPKernel) { // By default, restrict the maximum size to a value specified by // --gpu-max-threads-per-block=n or its default value for HIP. @@ -349,24 +330,8 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( F->addFnAttr("amdgpu-flat-work-group-size", AttrVal); } - if (const auto *Attr = FD->getAttr()) { - unsigned Min = - Attr->getMin()->EvaluateKnownConstInt(M.getContext()).getExtValue(); - unsigned Max = Attr->getMax() ? Attr->getMax() - ->EvaluateKnownConstInt(M.getContext()) - .getExtValue() - : 0; - - if (Min != 0) { - assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max"); - - std::string AttrVal = llvm::utostr(Min); - if (Max != 0) - AttrVal = AttrVal + "," + llvm::utostr(Max); - F->addFnAttr("amdgpu-waves-per-eu", AttrVal); - } else - assert(Max == 0 && "Max must be zero"); - } + if (const auto *Attr = FD->getAttr()) + M.handleAMDGPUWavesPerEUAttr(F, Attr); if (const auto *Attr = FD->getAttr()) { unsigned NumSGPR = Attr->getNumSGPR(); @@ -595,6 +560,47 @@ llvm::Value *AMDGPUTargetCodeGenInfo::createEnqueuedBlockKernel( return F; } +void CodeGenModule::handleAMDGPUFlatWorkGroupSizeAttr( + llvm::Function *F, const AMDGPUFlatWorkGroupSizeAttr *FlatWGS, + const ReqdWorkGroupSizeAttr *ReqdWGS) { + unsigned Min = 0; + unsigned Max = 0; + if (FlatWGS) { + Min = FlatWGS->getMin()->EvaluateKnownConstInt(getContext()).getExtValue(); + Max = FlatWGS->getMax()->EvaluateKnownConstInt(getContext()).getExtValue(); + } + if (ReqdWGS && Min == 0 && Max == 0) + Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim(); + + if (Min != 0) { + assert(Min <= Max && "Min must be less than or equal Max"); + + std::string AttrVal = llvm::utostr(Min) + "," + llvm::utostr(Max); + F->addFnAttr("amdgpu-flat-work-group-size", AttrVal); + } else + assert(Max == 0 && "Max must be zero"); +} + +void CodeGenModule::handleAMDGPUWavesPerEUAttr( + llvm::Function *F, const AMDGPUWavesPerEUAttr *Attr) { + unsigned Min = + Attr->getMin()->EvaluateKnownConstInt(getContext()).getExtValue(); + unsigned Max = + Attr->getMax() + ? Attr->getMax()->EvaluateKnownConstInt(getContext()).getExtValue() + : 0; + + if (Min != 0) { + assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max"); + + std::string AttrVal = llvm::utostr(Min); + if (Max != 0) + AttrVal = AttrVal + "," + llvm::utostr(Max); + F->addFnAttr("amdgpu-waves-per-eu", AttrVal); + } else + assert(Max == 0 && "Max must be zero"); +} + std::unique_ptr CodeGen::createAMDGPUTargetCodeGenInfo(CodeGenModule &CGM) { return std::make_unique(CGM.getTypes()); diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index 1ca0192333a0..0d4bbd795648 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -71,12 +71,12 @@ public: return true; } -private: // Adds a NamedMDNode with GV, Name, and Operand as operands, and adds the // resulting MDNode to the nvvm.annotations MDNode. static void addNVVMMetadata(llvm::GlobalValue *GV, StringRef Name, int Operand); +private: static void emitBuiltinSurfTexDeviceCopy(CodeGenFunction &CGF, LValue Dst, LValue Src) { llvm::Value *Handle = nullptr; @@ -256,24 +256,8 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes( // Create !{, metadata !"kernel", i32 1} node addNVVMMetadata(F, "kernel", 1); } - if (CUDALaunchBoundsAttr *Attr = FD->getAttr()) { - // Create !{, metadata !"maxntidx", i32 } node - llvm::APSInt MaxThreads(32); - MaxThreads = Attr->getMaxThreads()->EvaluateKnownConstInt(M.getContext()); - if (MaxThreads > 0) - addNVVMMetadata(F, "maxntidx", MaxThreads.getExtValue()); - - // min blocks is an optional argument for CUDALaunchBoundsAttr. If it was - // not specified in __launch_bounds__ or if the user specified a 0 value, - // we don't have to add a PTX directive. - if (Attr->getMinBlocks()) { - llvm::APSInt MinBlocks(32); - MinBlocks = Attr->getMinBlocks()->EvaluateKnownConstInt(M.getContext()); - if (MinBlocks > 0) - // Create !{, metadata !"minctasm", i32 } node - addNVVMMetadata(F, "minctasm", MinBlocks.getExtValue()); - } - } + if (CUDALaunchBoundsAttr *Attr = FD->getAttr()) + M.handleCUDALaunchBoundsAttr(F, Attr); } // Attach kernel metadata directly if compiling for NVPTX. @@ -303,6 +287,28 @@ bool NVPTXTargetCodeGenInfo::shouldEmitStaticExternCAliases() const { } } +void CodeGenModule::handleCUDALaunchBoundsAttr( + llvm::Function *F, const CUDALaunchBoundsAttr *Attr) { + // Create !{, metadata !"maxntidx", i32 } node + llvm::APSInt MaxThreads(32); + MaxThreads = Attr->getMaxThreads()->EvaluateKnownConstInt(getContext()); + if (MaxThreads > 0) + NVPTXTargetCodeGenInfo::addNVVMMetadata(F, "maxntidx", + MaxThreads.getExtValue()); + + // min blocks is an optional argument for CUDALaunchBoundsAttr. If it was + // not specified in __launch_bounds__ or if the user specified a 0 value, + // we don't have to add a PTX directive. + if (Attr->getMinBlocks()) { + llvm::APSInt MinBlocks(32); + MinBlocks = Attr->getMinBlocks()->EvaluateKnownConstInt(getContext()); + if (MinBlocks > 0) + // Create !{, metadata !"minctasm", i32 } node + NVPTXTargetCodeGenInfo::addNVVMMetadata(F, "minctasm", + MinBlocks.getExtValue()); + } +} + std::unique_ptr CodeGen::createNVPTXTargetCodeGenInfo(CodeGenModule &CGM) { return std::make_unique(CGM.getTypes()); diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 96d2e2cede62..66cabb194233 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -3411,6 +3411,9 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind, << getOpenMPClauseName(CKind) << getOpenMPDirectiveName(DKind); SkipUntil(tok::comma, tok::annot_pragma_openmp_end, StopBeforeMatch); break; + case OMPC_ompx_attribute: + Clause = ParseOpenMPOMPXAttributesClause(WrongDirective); + break; default: break; } @@ -3691,6 +3694,62 @@ OMPClause *Parser::ParseOpenMPInteropClause(OpenMPClauseKind Kind, llvm_unreachable("Unexpected interop variable clause."); } +OMPClause *Parser::ParseOpenMPOMPXAttributesClause(bool ParseOnly) { + SourceLocation Loc = ConsumeToken(); + // Parse '('. + BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end); + if (T.expectAndConsume(diag::err_expected_lparen_after, + getOpenMPClauseName(OMPC_ompx_attribute).data())) + return nullptr; + + ParsedAttributes ParsedAttrs(AttrFactory); + ParseAttributes(PAKM_GNU | PAKM_CXX11, ParsedAttrs); + + // Parse ')'. + if (T.consumeClose()) + return nullptr; + + if (ParseOnly) + return nullptr; + + SmallVector Attrs; + for (const ParsedAttr &PA : ParsedAttrs) { + switch (PA.getKind()) { + case ParsedAttr::AT_AMDGPUFlatWorkGroupSize: + if (!PA.checkExactlyNumArgs(Actions, 2)) + continue; + if (auto *A = Actions.CreateAMDGPUFlatWorkGroupSizeAttr( + PA, PA.getArgAsExpr(0), PA.getArgAsExpr(1))) + Attrs.push_back(A); + continue; + case ParsedAttr::AT_AMDGPUWavesPerEU: + if (!PA.checkAtLeastNumArgs(Actions, 1) || + !PA.checkAtMostNumArgs(Actions, 2)) + continue; + if (auto *A = Actions.CreateAMDGPUWavesPerEUAttr( + PA, PA.getArgAsExpr(0), + PA.getNumArgs() > 1 ? PA.getArgAsExpr(1) : nullptr)) + Attrs.push_back(A); + continue; + case ParsedAttr::AT_CUDALaunchBounds: + if (!PA.checkAtLeastNumArgs(Actions, 1) || + !PA.checkAtMostNumArgs(Actions, 2)) + continue; + if (auto *A = Actions.CreateLaunchBoundsAttr( + PA, PA.getArgAsExpr(0), + PA.getNumArgs() > 1 ? PA.getArgAsExpr(1) : nullptr)) + Attrs.push_back(A); + continue; + default: + Diag(Loc, diag::warn_omp_invalid_attribute_for_ompx_attributes) << PA; + continue; + }; + } + + return Actions.ActOnOpenMPXAttributeClause(Attrs, Loc, T.getOpenLocation(), + T.getCloseLocation()); +} + /// Parsing of simple OpenMP clauses like 'default' or 'proc_bind'. /// /// default-clause: diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index ed69e802c95d..429fa12ff2e2 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5633,21 +5633,28 @@ static Expr *makeLaunchBoundsArgExpr(Sema &S, Expr *E, return ValArg.getAs(); } -void Sema::AddLaunchBoundsAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *MaxThreads, Expr *MinBlocks) { +CUDALaunchBoundsAttr * +Sema::CreateLaunchBoundsAttr(const AttributeCommonInfo &CI, Expr *MaxThreads, + Expr *MinBlocks) { CUDALaunchBoundsAttr TmpAttr(Context, CI, MaxThreads, MinBlocks); MaxThreads = makeLaunchBoundsArgExpr(*this, MaxThreads, TmpAttr, 0); if (MaxThreads == nullptr) - return; + return nullptr; if (MinBlocks) { MinBlocks = makeLaunchBoundsArgExpr(*this, MinBlocks, TmpAttr, 1); if (MinBlocks == nullptr) - return; + return nullptr; } - D->addAttr(::new (Context) - CUDALaunchBoundsAttr(Context, CI, MaxThreads, MinBlocks)); + return ::new (Context) + CUDALaunchBoundsAttr(Context, CI, MaxThreads, MinBlocks); +} + +void Sema::AddLaunchBoundsAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *MaxThreads, Expr *MinBlocks) { + if (auto *Attr = CreateLaunchBoundsAttr(CI, MaxThreads, MinBlocks)) + D->addAttr(Attr); } static void handleLaunchBoundsAttr(Sema &S, Decl *D, const ParsedAttr &AL) { @@ -7862,16 +7869,22 @@ checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, return false; } -void Sema::addAMDGPUFlatWorkGroupSizeAttr(Decl *D, - const AttributeCommonInfo &CI, - Expr *MinExpr, Expr *MaxExpr) { +AMDGPUFlatWorkGroupSizeAttr * +Sema::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, + Expr *MinExpr, Expr *MaxExpr) { AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr); if (checkAMDGPUFlatWorkGroupSizeArguments(*this, MinExpr, MaxExpr, TmpAttr)) - return; + return nullptr; + return ::new (Context) + AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr); +} - D->addAttr(::new (Context) - AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr)); +void Sema::addAMDGPUFlatWorkGroupSizeAttr(Decl *D, + const AttributeCommonInfo &CI, + Expr *MinExpr, Expr *MaxExpr) { + if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr)) + D->addAttr(Attr); } static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D, @@ -7916,15 +7929,21 @@ static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr, return false; } -void Sema::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *MinExpr, Expr *MaxExpr) { +AMDGPUWavesPerEUAttr * +Sema::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI, Expr *MinExpr, + Expr *MaxExpr) { AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr); if (checkAMDGPUWavesPerEUArguments(*this, MinExpr, MaxExpr, TmpAttr)) - return; + return nullptr; - D->addAttr(::new (Context) - AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr)); + return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr); +} + +void Sema::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *MinExpr, Expr *MaxExpr) { + if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr)) + D->addAttr(Attr); } static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const ParsedAttr &AL) { diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index cf805987b378..3954bf2ad280 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -30,6 +30,7 @@ #include "clang/Sema/EnterExpressionEvaluationContext.h" #include "clang/Sema/Initialization.h" #include "clang/Sema/Lookup.h" +#include "clang/Sema/ParsedAttr.h" #include "clang/Sema/Scope.h" #include "clang/Sema/ScopeInfo.h" #include "clang/Sema/SemaInternal.h" @@ -24093,3 +24094,10 @@ OMPClause *Sema::ActOnOpenMPDoacrossClause( DSAStack->addDoacrossDependClause(C, OpsOffs); return C; } + +OMPClause *Sema::ActOnOpenMPXAttributeClause(ArrayRef Attrs, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + return new (Context) OMPXAttributeClause(Attrs, StartLoc, LParenLoc, EndLoc); +} diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 10b3587885e3..a73b54b668a4 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -2377,6 +2377,18 @@ public: EndLoc); } + /// Build a new OpenMP 'ompx_attribute' clause. + /// + /// By default, performs semantic analysis to build the new OpenMP clause. + /// Subclasses may override this routine to provide different behavior. + OMPClause *RebuildOMPXAttributeClause(ArrayRef Attrs, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + return getSema().ActOnOpenMPXAttributeClause(Attrs, StartLoc, LParenLoc, + EndLoc); + } + /// Build a new OpenMP 'align' clause. /// /// By default, performs semantic analysis to build the new OpenMP clause. @@ -10756,6 +10768,16 @@ TreeTransform::TransformOMPDoacrossClause(OMPDoacrossClause *C) { C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); } +template +OMPClause * +TreeTransform::TransformOMPXAttributeClause(OMPXAttributeClause *C) { + SmallVector NewAttrs; + for (auto *A : C->getAttrs()) + NewAttrs.push_back(getDerived().TransformAttr(A)); + return getDerived().RebuildOMPXAttributeClause( + NewAttrs, C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); +} + //===----------------------------------------------------------------------===// // Expression transformation //===----------------------------------------------------------------------===// diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 5f756961c6e1..dcb845dd551e 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -10370,6 +10370,9 @@ OMPClause *OMPClauseReader::readClause() { C = OMPDoacrossClause::CreateEmpty(Context, NumVars, NumLoops); break; } + case llvm::omp::OMPC_ompx_attribute: + C = new (Context) OMPXAttributeClause(); + break; #define OMP_CLAUSE_NO_CLASS(Enum, Str) \ case llvm::omp::Enum: \ break; @@ -11462,6 +11465,15 @@ void OMPClauseReader::VisitOMPDoacrossClause(OMPDoacrossClause *C) { C->setLoopData(I, Record.readSubExpr()); } +void OMPClauseReader::VisitOMPXAttributeClause(OMPXAttributeClause *C) { + AttrVec Attrs; + Record.readAttributes(Attrs); + C->setAttrs(Attrs); + C->setLocStart(Record.readSourceLocation()); + C->setLParenLoc(Record.readSourceLocation()); + C->setLocEnd(Record.readSourceLocation()); +} + OMPTraitInfo *ASTRecordReader::readOMPTraitInfo() { OMPTraitInfo &TI = getContext().getNewOMPTraitInfo(); TI.Sets.resize(readUInt32()); diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 26279d399b53..e238ad3d186f 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -7171,6 +7171,13 @@ void OMPClauseWriter::VisitOMPDoacrossClause(OMPDoacrossClause *C) { Record.AddStmt(C->getLoopData(I)); } +void OMPClauseWriter::VisitOMPXAttributeClause(OMPXAttributeClause *C) { + Record.AddAttributes(C->getAttrs()); + Record.AddSourceLocation(C->getBeginLoc()); + Record.AddSourceLocation(C->getLParenLoc()); + Record.AddSourceLocation(C->getEndLoc()); +} + void ASTRecordWriter::writeOMPTraitInfo(const OMPTraitInfo *TI) { writeUInt32(TI->Sets.size()); for (const auto &Set : TI->Sets) { diff --git a/clang/test/OpenMP/ompx_attributes_codegen.cpp b/clang/test/OpenMP/ompx_attributes_codegen.cpp new file mode 100644 index 000000000000..21e9805cbe82 --- /dev/null +++ b/clang/test/OpenMP/ompx_attributes_codegen.cpp @@ -0,0 +1,31 @@ +// REQUIRES: amdgpu-registered-target + +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s +// RUN: %clang_cc1 -target-cpu gfx900 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s +// expected-no-diagnostics + + +// Check that the target attributes are set on the generated kernel +void func() { + // CHECK: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l15() #0 + // CHECK: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l17() + // CHECK: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l19() #4 + + #pragma omp target ompx_attribute([[clang::amdgpu_flat_work_group_size(10, 20)]]) + {} + #pragma omp target teams ompx_attribute(__attribute__((launch_bounds(45, 90)))) + {} + #pragma omp target teams distribute parallel for simd ompx_attribute([[clang::amdgpu_flat_work_group_size(3, 17)]]) device(3) ompx_attribute(__attribute__((amdgpu_waves_per_eu(3, 7)))) + for (int i = 0; i < 1000; ++i) + {} +} + +// CHECK: attributes #0 +// CHECK-SAME: "amdgpu-flat-work-group-size"="10,20" +// CHECK: attributes #4 +// CHECK-SAME: "amdgpu-flat-work-group-size"="3,17" +// CHECK-SAME: "amdgpu-waves-per-eu"="3,7" + +// CHECK: !{ptr @__omp_offloading[[HASH]]_l17, !"maxntidx", i32 45} +// CHECK: !{ptr @__omp_offloading[[HASH]]_l17, !"minctasm", i32 90} diff --git a/clang/test/OpenMP/ompx_attributes_messages.cpp b/clang/test/OpenMP/ompx_attributes_messages.cpp new file mode 100644 index 000000000000..c59c19027d26 --- /dev/null +++ b/clang/test/OpenMP/ompx_attributes_messages.cpp @@ -0,0 +1,47 @@ +// RUN: %clang_cc1 -verify=expected -fopenmp -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized + +void bad() { + #pragma omp target data ompx_attribute() // expected-error {{unexpected OpenMP clause 'ompx_attribute' in directive '#pragma omp target data'}} + #pragma omp target data ompx_attribute(__attribute__((launch_bounds(1, 2)))) // expected-error {{unexpected OpenMP clause 'ompx_attribute' in directive '#pragma omp target data'}} expected-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}} + + #pragma omp target ompx_attribute() + {} + #pragma omp target ompx_attribute(__attribute__(())) + {} + #pragma omp target ompx_attribute(__attribute__((pure))) // expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'pure' is ignored}} + {} + #pragma omp target ompx_attribute(__attribute__((pure,amdgpu_waves_per_eu(1, 2), const))) // expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'pure' is ignored}} expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'const' is ignored}} + {} + #pragma omp target ompx_attribute(__attribute__((amdgpu_waves_per_eu()))) // expected-error {{'amdgpu_waves_per_eu' attribute takes at least 1 argument}} + {} + #pragma omp target ompx_attribute(__attribute__((amdgpu_waves_per_eu(1, 2, 3)))) // expected-error {{'amdgpu_waves_per_eu' attribute takes no more than 2 arguments}} + {} + #pragma omp target ompx_attribute(__attribute__((amdgpu_flat_work_group_size(1)))) // expected-error {{'amdgpu_flat_work_group_size' attribute requires exactly 2 arguments}} + {} + #pragma omp target ompx_attribute(__attribute__((amdgpu_flat_work_group_size(1, 2, 3,)))) // expected-error {{expected expression}} + {} + #pragma omp target ompx_attribute([[clang::amdgpu_waves_per_eu(1, 2, 3)]]) // expected-error {{'amdgpu_waves_per_eu' attribute takes no more than 2 arguments}} + {} + #pragma omp target ompx_attribute([[clang::unknown]]) // expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'unknown' is ignored}} + {} + #pragma omp target ompx_attribute(baz) // expected-error {{expected ')'}} expected-note {{to match this '('}} + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(1)))) + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(bad)))) // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}} + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2 // expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2) // expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2)) // expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2))) // expected-error {{expected ')'}} expected-note {{to match this '('}} + {} + #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, -3)))) // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}} + {} + #pragma omp target ompx_attribute(__attribute__((amdgpu_waves_per_eu(10, 1)))) // expected-error {{'amdgpu_waves_per_eu' attribute argument is invalid: min must not be greater than max}} + {} +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 0c629fe336fa..120ad4ea39e1 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2720,6 +2720,8 @@ void OMPClauseEnqueue::VisitOMPXDynCGroupMemClause( void OMPClauseEnqueue::VisitOMPDoacrossClause(const OMPDoacrossClause *C) { VisitOMPClauseList(C); } +void OMPClauseEnqueue::VisitOMPXAttributeClause(const OMPXAttributeClause *C) { +} } // namespace diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td index c67b54acc47c..68f7eca4daff 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMP.td +++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td @@ -448,6 +448,10 @@ def OMPC_Doacross : Clause<"doacross"> { let clangClass = "OMPDoacrossClause"; } +def OMPC_OMPX_Attribute : Clause<"ompx_attribute"> { + let clangClass = "OMPXAttributeClause"; +} + //===----------------------------------------------------------------------===// // Definition of OpenMP directives //===----------------------------------------------------------------------===// @@ -460,7 +464,8 @@ def OMP_Parallel : Directive<"parallel"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -645,7 +650,8 @@ def OMP_Target : Directive<"target"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -661,7 +667,8 @@ def OMP_Teams : Directive<"teams"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -744,7 +751,8 @@ def OMP_TargetParallel : Directive<"target parallel"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -779,7 +787,8 @@ def OMP_TargetParallelFor : Directive<"target parallel for"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -844,7 +853,8 @@ def OMP_ParallelFor : Directive<"parallel for"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelDo : Directive<"parallel do"> { @@ -889,7 +899,8 @@ def OMP_ParallelForSimd : Directive<"parallel for simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelDoSimd : Directive<"parallel do simd"> { @@ -929,7 +940,8 @@ def OMP_ParallelMaster : Directive<"parallel master"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelMasked : Directive<"parallel masked"> { @@ -944,7 +956,8 @@ def OMP_ParallelMasked : Directive<"parallel masked"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelSections : Directive<"parallel sections"> { @@ -958,7 +971,8 @@ def OMP_ParallelSections : Directive<"parallel sections"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause @@ -1127,7 +1141,8 @@ def OMP_DistributeParallelFor : Directive<"distribute parallel for"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_DistributeParallelDo : Directive<"distribute parallel do"> { @@ -1174,7 +1189,8 @@ def OMP_DistributeParallelForSimd : Directive<"distribute parallel for simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_DistributeParallelDoSimd : Directive<"distribute parallel do simd"> { @@ -1256,7 +1272,8 @@ def OMP_TargetParallelForSimd : Directive<"target parallel for simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1309,7 +1326,8 @@ def OMP_TargetSimd : Directive<"target simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1337,7 +1355,8 @@ def OMP_TeamsDistribute : Directive<"teams distribute"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_TeamsDistributeSimd : Directive<"teams distribute simd"> { @@ -1350,7 +1369,8 @@ def OMP_TeamsDistributeSimd : Directive<"teams distribute simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1388,7 +1408,8 @@ def OMP_TeamsDistributeParallelForSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_TeamsDistributeParallelDoSimd : @@ -1438,7 +1459,8 @@ def OMP_TeamsDistributeParallelFor : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_TeamsDistributeParallelDo : @@ -1479,7 +1501,8 @@ def OMP_TargetTeams : Directive<"target teams"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ @@ -1505,7 +1528,8 @@ def OMP_TargetTeamsDistribute : Directive<"target teams distribute"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1546,7 +1570,8 @@ def OMP_TargetTeamsDistributeParallelFor : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1617,7 +1642,8 @@ def OMP_TargetTeamsDistributeParallelForSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1678,7 +1704,8 @@ def OMP_TargetTeamsDistributeSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1773,7 +1800,8 @@ def OMP_ParallelMasterTaskloop : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelMaskedTaskloop : @@ -1798,7 +1826,8 @@ def OMP_ParallelMaskedTaskloop : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_MasterTaskloopSimd : Directive<"master taskloop simd"> { @@ -1883,7 +1912,8 @@ def OMP_ParallelMasterTaskloopSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelMaskedTaskloopSimd : @@ -1914,7 +1944,8 @@ def OMP_ParallelMaskedTaskloopSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_Depobj : Directive<"depobj"> { @@ -2021,6 +2052,7 @@ def OMP_teams_loop : Directive<"teams loop"> { VersionedClause, VersionedClause, VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -2045,7 +2077,8 @@ def OMP_target_teams_loop : Directive<"target teams loop"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -2068,6 +2101,7 @@ def OMP_parallel_loop : Directive<"parallel loop"> { VersionedClause, VersionedClause, VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -2094,6 +2128,7 @@ def OMP_target_parallel_loop : Directive<"target parallel loop"> { VersionedClause, VersionedClause, VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause,