shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[OpenMP] Add the ompx_attribute clause for target directives

Browse Source 
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
This commit is contained in:
Johannes Doerfert 2023-03-02 18:35:15 -08:00
parent fb2a971c01
commit ef9ec4bbcc
22 changed files with 489 additions and 105 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

48
clang/include/clang/AST/OpenMPClause.h
View File

@ -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<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

6
clang/include/clang/AST/RecursiveASTVisitor.h
View File

@ -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

3
clang/include/clang/Basic/DiagnosticGroups.td
View File

@ -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.

3
clang/include/clang/Basic/DiagnosticParseKinds.td
View File

@ -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<

7
clang/include/clang/Parse/Parser.h
View File

@ -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.

21
clang/include/clang/Sema/Sema.h
View File

@ -10988,6 +10988,11 @@ public:
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 @@ 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<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.

12
clang/lib/AST/OpenMPClause.cpp
View File

@ -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) {

2
clang/lib/AST/StmtProfile.cpp
View File

@ -928,6 +928,8 @@ void OMPClauseProfiler::VisitOMPXDynCGroupMemClause(
void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
VisitOMPClauseList(C);
}
void OMPClauseProfiler::VisitOMPXAttributeClause(const OMPXAttributeClause *C) {
}
} // namespace
void

19
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -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

15
clang/lib/CodeGen/CodeGenModule.h
View File

@ -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,

82
clang/lib/CodeGen/Targets/AMDGPU.cpp
View File

@ -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());

44
clang/lib/CodeGen/Targets/NVPTX.cpp
View File

@ -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 !{<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());

59
clang/lib/Parse/ParseOpenMP.cpp
View File

@ -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:

53
clang/lib/Sema/SemaDeclAttr.cpp
View File

@ -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;
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) {

8
clang/lib/Sema/SemaOpenMP.cpp
View File

@ -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);
}

22
clang/lib/Sema/TreeTransform.h
View File

@ -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<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
//===----------------------------------------------------------------------===//

12
clang/lib/Serialization/ASTReader.cpp
View File

@ -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());

7
clang/lib/Serialization/ASTWriter.cpp
View File

@ -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) {

31
clang/test/OpenMP/ompx_attributes_codegen.cpp Normal file
View File

@ -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}

47
clang/test/OpenMP/ompx_attributes_messages.cpp Normal file
View File

@ -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}}
{}
}

2
clang/tools/libclang/CIndex.cpp
View File

@ -2720,6 +2720,8 @@ void OMPClauseEnqueue::VisitOMPXDynCGroupMemClause(
void OMPClauseEnqueue::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
VisitOMPClauseList(C);
}
void OMPClauseEnqueue::VisitOMPXAttributeClause(const OMPXAttributeClause *C) {
}
} // namespace

91
llvm/include/llvm/Frontend/OpenMP/OMP.td
View File

@ -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>,