From 24033980849b1fd1a6501c5f075f2551df91837d Mon Sep 17 00:00:00 2001 From: Johannes Doerfert Date: Tue, 25 Jul 2023 09:37:49 -0700 Subject: [PATCH] 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 --- .../clang/Basic/DiagnosticParseKinds.td | 3 + clang/include/clang/Parse/Parser.h | 7 ++ clang/include/clang/Sema/Sema.h | 10 +++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 41 +-------- clang/lib/CodeGen/Targets/NVPTX.cpp | 22 +---- clang/lib/Sema/SemaDeclAttr.cpp | 34 +++++--- clang/lib/Sema/SemaOpenMP.cpp | 1 + .../test/OpenMP/ompx_attributes_messages.cpp | 47 ++++++++++ llvm/include/llvm/Frontend/OpenMP/OMP.td | 87 +++++++++++++------ 9 files changed, 156 insertions(+), 96 deletions(-) create mode 100644 clang/test/OpenMP/ompx_attributes_messages.cpp diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td index b44950ce9f39..7d5d47189111 100644 --- a/clang/include/clang/Basic/DiagnosticParseKinds.td +++ b/clang/include/clang/Basic/DiagnosticParseKinds.td @@ -1549,6 +1549,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 b5ac769faf7e..f8e02a788703 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3516,6 +3516,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 74ab580b84ec..1d8903907da4 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -11157,11 +11157,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, diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 460d69607963..0411846cf9b0 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -327,26 +327,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. @@ -359,24 +340,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(); diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index d9eee5f3c196..d0dc7c258a03 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 !{, 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. diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 82b0d5787d84..842a01a88cd3 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7883,16 +7883,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, @@ -7937,15 +7943,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 e89c83fe9d23..1bd34f73e5f7 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" 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/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td index fa9abd986830..f8b3b0c75249 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMP.td +++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td @@ -466,7 +466,8 @@ def OMP_Parallel : Directive<"parallel"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -650,7 +651,8 @@ def OMP_Target : Directive<"target"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -667,7 +669,8 @@ def OMP_Teams : Directive<"teams"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -752,7 +755,8 @@ def OMP_TargetParallel : Directive<"target parallel"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -787,7 +791,8 @@ def OMP_TargetParallelFor : Directive<"target parallel for"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -853,7 +858,8 @@ def OMP_ParallelFor : Directive<"parallel for"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelDo : Directive<"parallel do"> { @@ -898,7 +904,8 @@ def OMP_ParallelForSimd : Directive<"parallel for simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelDoSimd : Directive<"parallel do simd"> { @@ -938,7 +945,8 @@ def OMP_ParallelMaster : Directive<"parallel master"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelMasked : Directive<"parallel masked"> { @@ -953,7 +961,8 @@ def OMP_ParallelMasked : Directive<"parallel masked"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelSections : Directive<"parallel sections"> { @@ -966,7 +975,8 @@ def OMP_ParallelSections : Directive<"parallel sections"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1137,7 +1147,8 @@ def OMP_DistributeParallelFor : Directive<"distribute parallel for"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_DistributeParallelDo : Directive<"distribute parallel do"> { @@ -1184,7 +1195,8 @@ def OMP_DistributeParallelForSimd : Directive<"distribute parallel for simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_DistributeParallelDoSimd : Directive<"distribute parallel do simd"> { @@ -1266,7 +1278,8 @@ def OMP_TargetParallelForSimd : Directive<"target parallel for simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1321,7 +1334,8 @@ def OMP_TargetSimd : Directive<"target simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1349,7 +1363,8 @@ def OMP_TeamsDistribute : Directive<"teams distribute"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause @@ -1366,7 +1381,8 @@ def OMP_TeamsDistributeSimd : Directive<"teams distribute simd"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1403,7 +1419,8 @@ def OMP_TeamsDistributeParallelForSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_TeamsDistributeParallelDoSimd : @@ -1453,7 +1470,8 @@ def OMP_TeamsDistributeParallelFor : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_TeamsDistributeParallelDo : @@ -1494,7 +1512,8 @@ def OMP_TargetTeams : Directive<"target teams"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ @@ -1521,7 +1540,8 @@ def OMP_TargetTeamsDistribute : Directive<"target teams distribute"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1562,7 +1582,8 @@ def OMP_TargetTeamsDistributeParallelFor : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1633,7 +1654,8 @@ def OMP_TargetTeamsDistributeParallelForSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1694,7 +1716,8 @@ def OMP_TargetTeamsDistributeSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -1794,7 +1817,8 @@ def OMP_ParallelMasterTaskloop : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelMaskedTaskloop : @@ -1819,7 +1843,8 @@ def OMP_ParallelMaskedTaskloop : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_MasterTaskloopSimd : Directive<"master taskloop simd"> { @@ -1904,7 +1929,8 @@ def OMP_ParallelMasterTaskloopSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_ParallelMaskedTaskloopSimd : @@ -1935,7 +1961,8 @@ def OMP_ParallelMaskedTaskloopSimd : VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; } def OMP_Depobj : Directive<"depobj"> { @@ -2051,6 +2078,7 @@ def OMP_teams_loop : Directive<"teams loop"> { VersionedClause, VersionedClause, VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -2076,7 +2104,8 @@ def OMP_target_teams_loop : Directive<"target teams loop"> { VersionedClause, VersionedClause, VersionedClause, - VersionedClause + VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -2098,6 +2127,7 @@ def OMP_parallel_loop : Directive<"parallel loop"> { VersionedClause, VersionedClause, VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause, @@ -2125,6 +2155,7 @@ def OMP_target_parallel_loop : Directive<"target parallel loop"> { VersionedClause, VersionedClause, VersionedClause, + VersionedClause, ]; let allowedOnceClauses = [ VersionedClause,