Author: Johannes Doerfert Date: 2023-07-25T10:40:35-07:00 New Revision: 08a220764b1e266c4694f614fd4fda7bd2122580
URL: https://github.com/llvm/llvm-project/commit/08a220764b1e266c4694f614fd4fda7bd2122580 DIFF: https://github.com/llvm/llvm-project/commit/08a220764b1e266c4694f614fd4fda7bd2122580.diff LOG: Reapply "[OpenMP] Add the `ompx_attribute` clause for target directives" This reverts commit 0d12683046ca75fb08e285f4622f2af5c82609dc and reapplies ef9ec4bbcca2fa4f64df47bc426f1d1c59ea47e2 with an extension to fix the Flang build. Differential Revision: https://reviews.llvm.org/D156184 Added: clang/test/OpenMP/ompx_attributes_codegen.cpp clang/test/OpenMP/ompx_attributes_messages.cpp Modified: clang/include/clang/AST/OpenMPClause.h clang/include/clang/AST/RecursiveASTVisitor.h clang/include/clang/Basic/DiagnosticGroups.td clang/include/clang/Basic/DiagnosticParseKinds.td clang/include/clang/Parse/Parser.h clang/include/clang/Sema/Sema.h clang/lib/AST/OpenMPClause.cpp clang/lib/AST/StmtProfile.cpp clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/lib/CodeGen/CodeGenModule.h clang/lib/CodeGen/Targets/AMDGPU.cpp clang/lib/CodeGen/Targets/NVPTX.cpp clang/lib/Parse/ParseOpenMP.cpp clang/lib/Sema/SemaDeclAttr.cpp clang/lib/Sema/SemaOpenMP.cpp clang/lib/Sema/TreeTransform.h clang/lib/Serialization/ASTReader.cpp clang/lib/Serialization/ASTWriter.cpp clang/tools/libclang/CIndex.cpp flang/lib/Semantics/check-omp-structure.cpp llvm/include/llvm/Frontend/OpenMP/OMP.td Removed: ################################################################################ diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h index 0bea21270692cf..31ae3d42e232fc 100644 --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -9172,6 +9172,54 @@ class OMPDoacrossClause final } }; +/// 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<llvm::omp::OMPC_ompx_attribute> { + friend class OMPClauseReader; + + /// Location of '('. + SourceLocation LParenLoc; + + /// The parsed attributes (clause arguments) + SmallVector<const Attr *> 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<const Attr *> 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<const Attr *> getAttrs() const { return Attrs; } + +private: + /// Replace the attributes with \p NewAttrs. + void setAttrs(ArrayRef<Attr *> 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 604875cd6337a4..fc2d1ff708bf7a 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -3875,6 +3875,12 @@ bool RecursiveASTVisitor<Derived>::VisitOMPDoacrossClause( return true; } +template <typename Derived> +bool RecursiveASTVisitor<Derived>::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 7b4d415bf06494..6a0a01e4a981a4 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 8d729c31641ed8..a804442ab34ec7 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<OpenMPClauses>; def warn_omp_depend_in_ordered_deprecated : Warning<"'depend' clause for" " 'ordered' is deprecated; use 'doacross' instead">, InGroup<Deprecated>; +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<OpenMPExtensions>; // 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 475dfe845528d9..b58041477c9ec9 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3490,6 +3490,13 @@ class Parser : public CodeCompletionHandler { // 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 3418a37b307785..7c641d5e273a03 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10988,6 +10988,11 @@ class Sema final { bool ConstantFoldAttrArgs(const AttributeCommonInfo &CI, MutableArrayRef<Expr *> 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 @@ class Sema final { 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 @@ class Sema final { ArrayRef<Expr *> VarList, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); + /// Called on a well-formed 'ompx_attribute' clause. + OMPClause *ActOnOpenMPXAttributeClause(ArrayRef<const Attr *> 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 4c895822ffdf85..f5ad75028a641e 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 d8a667b2d0fdc4..60646f7a0da57c 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 a52ec8909b12ac..a4f7eb96c0d9bc 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<OMPXAttributeClause>()) { + for (auto *A : C->getAttrs()) { + if (auto *Attr = dyn_cast<CUDALaunchBoundsAttr>(A)) + CGM.handleCUDALaunchBoundsAttr(OutlinedFn, Attr); + else if (auto *Attr = dyn_cast<AMDGPUFlatWorkGroupSizeAttr>(A)) + CGM.handleAMDGPUFlatWorkGroupSizeAttr(OutlinedFn, Attr); + else if (auto *Attr = dyn_cast<AMDGPUWavesPerEUAttr>(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 05cb217e2bee4e..f5fd94492540f2 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1557,6 +1557,21 @@ class CodeGenModule : public CodeGenTypeCache { /// 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 796a2be81a09c7..bac7787643e33a 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<AMDGPUFlatWorkGroupSizeAttr>(); 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<AMDGPUWavesPerEUAttr>()) { - 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<AMDGPUWavesPerEUAttr>()) + M.handleAMDGPUWavesPerEUAttr(F, Attr); if (const auto *Attr = FD->getAttr<AMDGPUNumSGPRAttr>()) { 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<TargetCodeGenInfo> CodeGen::createAMDGPUTargetCodeGenInfo(CodeGenModule &CGM) { return std::make_unique<AMDGPUTargetCodeGenInfo>(CGM.getTypes()); diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index 1ca0192333a0f8..0d4bbd79564800 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -71,12 +71,12 @@ class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo { 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 !{<func-ref>, metadata !"kernel", i32 1} node addNVVMMetadata(F, "kernel", 1); } - if (CUDALaunchBoundsAttr *Attr = FD->getAttr<CUDALaunchBoundsAttr>()) { - // Create !{<func-ref>, metadata !"maxntidx", i32 <val>} 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 !{<func-ref>, metadata !"minctasm", i32 <val>} node - addNVVMMetadata(F, "minctasm", MinBlocks.getExtValue()); - } - } + if (CUDALaunchBoundsAttr *Attr = FD->getAttr<CUDALaunchBoundsAttr>()) + 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 !{<func-ref>, metadata !"maxntidx", i32 <val>} 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 !{<func-ref>, metadata !"minctasm", i32 <val>} node + NVPTXTargetCodeGenInfo::addNVVMMetadata(F, "minctasm", + MinBlocks.getExtValue()); + } +} + std::unique_ptr<TargetCodeGenInfo> CodeGen::createNVPTXTargetCodeGenInfo(CodeGenModule &CGM) { return std::make_unique<NVPTXTargetCodeGenInfo>(CGM.getTypes()); diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 96d2e2cede6289..66cabb19423348 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<Attr *> 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 ed69e802c95dd5..429fa12ff2e293 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<Expr>(); } -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; + + return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr); +} - D->addAttr(::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 cf805987b378e3..3954bf2ad28041 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<const Attr *> 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 10b3587885e39f..a73b54b668a415 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -2377,6 +2377,18 @@ class TreeTransform { 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 diff erent behavior. + OMPClause *RebuildOMPXAttributeClause(ArrayRef<const Attr *> 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<Derived>::TransformOMPDoacrossClause(OMPDoacrossClause *C) { C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); } +template <typename Derived> +OMPClause * +TreeTransform<Derived>::TransformOMPXAttributeClause(OMPXAttributeClause *C) { + SmallVector<const Attr *> 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 5f756961c6e1d0..dcb845dd551e8d 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 26279d399b53a9..e238ad3d186fba 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 00000000000000..21e9805cbe8293 --- /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 00000000000000..c59c19027d26f8 --- /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 39886b23bb36f1..1bdc0bf742a8ce 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/flang/lib/Semantics/check-omp-structure.cpp b/flang/lib/Semantics/check-omp-structure.cpp index 7337102d40e8e7..b4b838a45f7eae 100644 --- a/flang/lib/Semantics/check-omp-structure.cpp +++ b/flang/lib/Semantics/check-omp-structure.cpp @@ -1962,6 +1962,7 @@ CHECK_SIMPLE_CLAUSE(Align, OMPC_align) CHECK_SIMPLE_CLAUSE(Compare, OMPC_compare) CHECK_SIMPLE_CLAUSE(CancellationConstructType, OMPC_cancellation_construct_type) CHECK_SIMPLE_CLAUSE(Doacross, OMPC_doacross) +CHECK_SIMPLE_CLAUSE(OmpxAttribute, OMPC_ompx_attribute) CHECK_REQ_SCALAR_INT_CLAUSE(Grainsize, OMPC_grainsize) CHECK_REQ_SCALAR_INT_CLAUSE(NumTasks, OMPC_num_tasks) diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td index c67b54acc47c0d..68f7eca4daffbf 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<OMPC_Shared>, VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Copyin>, - VersionedClause<OMPC_Allocate> + VersionedClause<OMPC_Allocate>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Default>, @@ -645,7 +650,8 @@ def OMP_Target : Directive<"target"> { VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_InReduction, 50>, VersionedClause<OMPC_Allocate>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Device>, @@ -661,7 +667,8 @@ def OMP_Teams : Directive<"teams"> { VersionedClause<OMPC_FirstPrivate>, VersionedClause<OMPC_Shared>, VersionedClause<OMPC_Reduction>, - VersionedClause<OMPC_Allocate> + VersionedClause<OMPC_Allocate>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Default>, @@ -744,7 +751,8 @@ def OMP_TargetParallel : Directive<"target parallel"> { VersionedClause<OMPC_IsDevicePtr>, VersionedClause<OMPC_HasDeviceAddr, 51>, VersionedClause<OMPC_Allocate>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_DefaultMap>, @@ -779,7 +787,8 @@ def OMP_TargetParallelFor : Directive<"target parallel for"> { VersionedClause<OMPC_HasDeviceAddr, 51>, VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_Order, 50>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, @@ -844,7 +853,8 @@ def OMP_ParallelFor : Directive<"parallel for"> { VersionedClause<OMPC_Ordered>, VersionedClause<OMPC_Linear>, VersionedClause<OMPC_Allocate>, - VersionedClause<OMPC_Order, 50> + VersionedClause<OMPC_Order, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_ParallelDo : Directive<"parallel do"> { @@ -889,7 +899,8 @@ def OMP_ParallelForSimd : Directive<"parallel for simd"> { VersionedClause<OMPC_Ordered>, VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_NonTemporal, 50>, - VersionedClause<OMPC_Order, 50> + VersionedClause<OMPC_Order, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_ParallelDoSimd : Directive<"parallel do simd"> { @@ -929,7 +940,8 @@ def OMP_ParallelMaster : Directive<"parallel master"> { VersionedClause<OMPC_Copyin>, VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_ProcBind>, - VersionedClause<OMPC_Allocate> + VersionedClause<OMPC_Allocate>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_ParallelMasked : Directive<"parallel masked"> { @@ -944,7 +956,8 @@ def OMP_ParallelMasked : Directive<"parallel masked"> { VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_ProcBind>, VersionedClause<OMPC_Allocate>, - VersionedClause<OMPC_Filter> + VersionedClause<OMPC_Filter>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_ParallelSections : Directive<"parallel sections"> { @@ -958,7 +971,8 @@ def OMP_ParallelSections : Directive<"parallel sections"> { VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Copyin>, VersionedClause<OMPC_LastPrivate>, - VersionedClause<OMPC_Allocate> + VersionedClause<OMPC_Allocate>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_NumThreads> @@ -1127,7 +1141,8 @@ def OMP_DistributeParallelFor : Directive<"distribute parallel for"> { VersionedClause<OMPC_Copyin>, VersionedClause<OMPC_Schedule>, VersionedClause<OMPC_Allocate>, - VersionedClause<OMPC_Order, 50> + VersionedClause<OMPC_Order, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_DistributeParallelDo : Directive<"distribute parallel do"> { @@ -1174,7 +1189,8 @@ def OMP_DistributeParallelForSimd : Directive<"distribute parallel for simd"> { VersionedClause<OMPC_SimdLen>, VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_NonTemporal, 50>, - VersionedClause<OMPC_Order, 50> + VersionedClause<OMPC_Order, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_DistributeParallelDoSimd : Directive<"distribute parallel do simd"> { @@ -1256,7 +1272,8 @@ def OMP_TargetParallelForSimd : Directive<"target parallel for simd"> { VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_NonTemporal, 50>, VersionedClause<OMPC_Order, 50>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, @@ -1309,7 +1326,8 @@ def OMP_TargetSimd : Directive<"target simd"> { VersionedClause<OMPC_Private>, VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Shared>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Collapse>, @@ -1337,7 +1355,8 @@ def OMP_TeamsDistribute : Directive<"teams distribute"> { VersionedClause<OMPC_LastPrivate>, VersionedClause<OMPC_Collapse>, VersionedClause<OMPC_DistSchedule>, - VersionedClause<OMPC_Allocate> + VersionedClause<OMPC_Allocate>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_TeamsDistributeSimd : Directive<"teams distribute simd"> { @@ -1350,7 +1369,8 @@ def OMP_TeamsDistributeSimd : Directive<"teams distribute simd"> { VersionedClause<OMPC_NonTemporal, 50>, VersionedClause<OMPC_Private>, VersionedClause<OMPC_Reduction>, - VersionedClause<OMPC_Shared> + VersionedClause<OMPC_Shared>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Collapse>, @@ -1388,7 +1408,8 @@ def OMP_TeamsDistributeParallelForSimd : VersionedClause<OMPC_ThreadLimit>, VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_NonTemporal, 50>, - VersionedClause<OMPC_Order, 50> + VersionedClause<OMPC_Order, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_TeamsDistributeParallelDoSimd : @@ -1438,7 +1459,8 @@ def OMP_TeamsDistributeParallelFor : VersionedClause<OMPC_ThreadLimit>, VersionedClause<OMPC_Copyin>, VersionedClause<OMPC_Allocate>, - VersionedClause<OMPC_Order, 50> + VersionedClause<OMPC_Order, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_TeamsDistributeParallelDo : @@ -1479,7 +1501,8 @@ def OMP_TargetTeams : Directive<"target teams"> { VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_UsesAllocators, 50>, - VersionedClause<OMPC_Shared> + VersionedClause<OMPC_Shared>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ @@ -1505,7 +1528,8 @@ def OMP_TargetTeamsDistribute : Directive<"target teams distribute"> { VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_UsesAllocators, 50>, VersionedClause<OMPC_Shared>, - VersionedClause<OMPC_LastPrivate> + VersionedClause<OMPC_LastPrivate>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Device>, @@ -1546,7 +1570,8 @@ def OMP_TargetTeamsDistributeParallelFor : VersionedClause<OMPC_Schedule>, VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_Order, 50>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, @@ -1617,7 +1642,8 @@ def OMP_TargetTeamsDistributeParallelForSimd : VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_NonTemporal, 50>, VersionedClause<OMPC_Order, 50>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, @@ -1678,7 +1704,8 @@ def OMP_TargetTeamsDistributeSimd : VersionedClause<OMPC_Private>, VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Shared>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Device>, @@ -1773,7 +1800,8 @@ def OMP_ParallelMasterTaskloop : VersionedClause<OMPC_Allocate>, VersionedClause<OMPC_NumThreads>, VersionedClause<OMPC_ProcBind>, - VersionedClause<OMPC_Copyin> + VersionedClause<OMPC_Copyin>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_ParallelMaskedTaskloop : @@ -1798,7 +1826,8 @@ def OMP_ParallelMaskedTaskloop : VersionedClause<OMPC_NumThreads>, VersionedClause<OMPC_ProcBind>, VersionedClause<OMPC_Copyin>, - VersionedClause<OMPC_Filter> + VersionedClause<OMPC_Filter>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_MasterTaskloopSimd : Directive<"master taskloop simd"> { @@ -1883,7 +1912,8 @@ def OMP_ParallelMasterTaskloopSimd : VersionedClause<OMPC_SafeLen>, VersionedClause<OMPC_SimdLen>, VersionedClause<OMPC_NonTemporal, 50>, - VersionedClause<OMPC_Order, 50> + VersionedClause<OMPC_Order, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_ParallelMaskedTaskloopSimd : @@ -1914,7 +1944,8 @@ def OMP_ParallelMaskedTaskloopSimd : VersionedClause<OMPC_SimdLen>, VersionedClause<OMPC_NonTemporal, 50>, VersionedClause<OMPC_Order, 50>, - VersionedClause<OMPC_Filter> + VersionedClause<OMPC_Filter>, + VersionedClause<OMPC_OMPX_Attribute>, ]; } def OMP_Depobj : Directive<"depobj"> { @@ -2021,6 +2052,7 @@ def OMP_teams_loop : Directive<"teams loop"> { VersionedClause<OMPC_Private>, VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Shared>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Bind, 50>, @@ -2045,7 +2077,8 @@ def OMP_target_teams_loop : Directive<"target teams loop"> { VersionedClause<OMPC_Private>, VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Shared>, - VersionedClause<OMPC_UsesAllocators, 50> + VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Bind, 50>, @@ -2068,6 +2101,7 @@ def OMP_parallel_loop : Directive<"parallel loop"> { VersionedClause<OMPC_Private>, VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Shared>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Bind, 50>, @@ -2094,6 +2128,7 @@ def OMP_target_parallel_loop : Directive<"target parallel loop"> { VersionedClause<OMPC_Reduction>, VersionedClause<OMPC_Shared>, VersionedClause<OMPC_UsesAllocators, 50>, + VersionedClause<OMPC_OMPX_Attribute>, ]; let allowedOnceClauses = [ VersionedClause<OMPC_Bind, 50>, _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
