maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OPENMP50]Add device/kind context selector support. 

Summary: Added basic parsing/sema support for device/kind context selector.

Reviewers: jdoerfert

Subscribers: rampitec, aheejin, fedor.sergeev, simoncook, guansong, s.egerton, hfinkel, kkwli0, caomhin, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D70245
This commit is contained in:
Alexey Bataev 2019-11-05 15:13:30 -05:00
parent 70d173fb1f
commit 4e8231b5cf
14 changed files with 905 additions and 74 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

26
clang/include/clang/Basic/Attr.td
View File

@ -3337,7 +3337,8 @@ def OMPDeclareVariant : InheritableAttr {
VariadicExprArgument<"Scores">, VariadicExprArgument<"Scores">,
VariadicUnsignedArgument<"CtxSelectorSets">, VariadicUnsignedArgument<"CtxSelectorSets">,
VariadicUnsignedArgument<"CtxSelectors">, VariadicUnsignedArgument<"CtxSelectors">,
VariadicStringArgument<"ImplVendors"> VariadicStringArgument<"ImplVendors">,
VariadicStringArgument<"DeviceKinds">
]; ];
let AdditionalMembers = [{ let AdditionalMembers = [{
void printScore(raw_ostream & OS, const PrintingPolicy &Policy, unsigned I) const { void printScore(raw_ostream & OS, const PrintingPolicy &Policy, unsigned I) const {
@ -3377,6 +3378,27 @@ def OMPDeclareVariant : InheritableAttr {
} }
OS << ")"; OS << ")";
break; break;
case OMP_CTX_kind:
llvm_unreachable("Unexpected context selector in implementation set.");
case OMP_CTX_unknown:
llvm_unreachable("Unknown context selector.");
}
OS << "}";
break;
case OMP_CTX_SET_device:
OS << "device={";
switch (Ctx) {
case OMP_CTX_kind:
OS << "kind(";
if (deviceKinds_size() > 0) {
OS << *deviceKinds().begin();
for (StringRef KindName : llvm::drop_begin(deviceKinds(), 1))
OS << ", " << KindName;
}
OS << ")";
break;
case OMP_CTX_vendor:
llvm_unreachable("Unexpected context selector in device set.");
case OMP_CTX_unknown: case OMP_CTX_unknown:
llvm_unreachable("Unknown context selector."); llvm_unreachable("Unknown context selector.");
} }
@ -3385,6 +3407,8 @@ def OMPDeclareVariant : InheritableAttr {
case OMP_CTX_SET_unknown: case OMP_CTX_SET_unknown:
llvm_unreachable("Unknown context selector set."); llvm_unreachable("Unknown context selector set.");
} }
if (I != E - 1)
OS << ",";
} }
OS << ")"; OS << ")";
} }

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

@ -1225,6 +1225,9 @@ def note_omp_declare_variant_ctx_used_here : Note<
def warn_omp_more_one_device_type_clause : Warning< def warn_omp_more_one_device_type_clause : Warning<
"more than one 'device_type' clause is specified">, "more than one 'device_type' clause is specified">,
InGroup<OpenMPClauses>; InGroup<OpenMPClauses>;
def err_omp_wrong_device_kind_trait : Error<
"unknown '%0' device kind trait in the 'device' context selector set, expected"
" one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'">;
// Pragma loop support. // Pragma loop support.
def err_pragma_loop_missing_argument : Error< def err_pragma_loop_missing_argument : Error<

2
clang/include/clang/Basic/OpenMPKinds.def
View File

@ -221,9 +221,11 @@
// OpenMP context selector sets. // OpenMP context selector sets.
OPENMP_CONTEXT_SELECTOR_SET(implementation) OPENMP_CONTEXT_SELECTOR_SET(implementation)
OPENMP_CONTEXT_SELECTOR_SET(device)
// OpenMP context selectors. // OpenMP context selectors.
OPENMP_CONTEXT_SELECTOR(vendor) OPENMP_CONTEXT_SELECTOR(vendor)
OPENMP_CONTEXT_SELECTOR(kind)
// OpenMP directives. // OpenMP directives.
OPENMP_DIRECTIVE(threadprivate) OPENMP_DIRECTIVE(threadprivate)

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

@ -20,6 +20,7 @@
#include "clang/AST/StmtOpenMP.h" #include "clang/AST/StmtOpenMP.h"
#include "clang/Basic/BitmaskEnum.h" #include "clang/Basic/BitmaskEnum.h"
#include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/SetOperations.h"
#include "llvm/Bitcode/BitcodeReader.h" #include "llvm/Bitcode/BitcodeReader.h"
#include "llvm/IR/DerivedTypes.h" #include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/GlobalValue.h" #include "llvm/IR/GlobalValue.h"
@ -11032,8 +11033,10 @@ using CompleteOMPContextSelectorData = SmallVector<OMPContextSelectorData, 4>;
} // anonymous namespace } // anonymous namespace
/// Checks current context and returns true if it matches the context selector. /// Checks current context and returns true if it matches the context selector.
template <OpenMPContextSelectorSetKind CtxSet, OpenMPContextSelectorKind Ctx> template <OpenMPContextSelectorSetKind CtxSet, OpenMPContextSelectorKind Ctx,
static bool checkContext(const OMPContextSelectorData &Data) { typename... Arguments>
static bool checkContext(const OMPContextSelectorData &Data,
Arguments... Params) {
assert(Data.CtxSet != OMP_CTX_SET_unknown && Data.Ctx != OMP_CTX_unknown && assert(Data.CtxSet != OMP_CTX_SET_unknown && Data.Ctx != OMP_CTX_unknown &&
"Unknown context selector or context selector set."); "Unknown context selector or context selector set.");
return false; return false;
@ -11048,7 +11051,92 @@ bool checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(
[](StringRef S) { return !S.compare_lower("llvm"); }); [](StringRef S) { return !S.compare_lower("llvm"); });
} }
bool matchesContext(const CompleteOMPContextSelectorData &ContextData) { /// Checks for device={kind(<kind>)} context selector.
/// \returns true if <kind>="host" and compilation is for host.
/// true if <kind>="nohost" and compilation is for device.
/// true if <kind>="cpu" and compilation is for Arm, X86 or PPC CPU.
/// true if <kind>="gpu" and compilation is for NVPTX or AMDGCN.
/// false otherwise.
template <>
bool checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(
const OMPContextSelectorData &Data, CodeGenModule &CGM) {
for (StringRef Name : Data.Names) {
if (!Name.compare_lower("host")) {
if (CGM.getLangOpts().OpenMPIsDevice)
return false;
continue;
}
if (!Name.compare_lower("nohost")) {
if (!CGM.getLangOpts().OpenMPIsDevice)
return false;
continue;
}
switch (CGM.getTriple().getArch()) {
case llvm::Triple::arm:
case llvm::Triple::armeb:
case llvm::Triple::aarch64:
case llvm::Triple::aarch64_be:
case llvm::Triple::aarch64_32:
case llvm::Triple::ppc:
case llvm::Triple::ppc64:
case llvm::Triple::ppc64le:
case llvm::Triple::x86:
case llvm::Triple::x86_64:
if (Name.compare_lower("cpu"))
return false;
break;
case llvm::Triple::amdgcn:
case llvm::Triple::nvptx:
case llvm::Triple::nvptx64:
if (Name.compare_lower("gpu"))
return false;
break;
case llvm::Triple::UnknownArch:
case llvm::Triple::arc:
case llvm::Triple::avr:
case llvm::Triple::bpfel:
case llvm::Triple::bpfeb:
case llvm::Triple::hexagon:
case llvm::Triple::mips:
case llvm::Triple::mipsel:
case llvm::Triple::mips64:
case llvm::Triple::mips64el:
case llvm::Triple::msp430:
case llvm::Triple::r600:
case llvm::Triple::riscv32:
case llvm::Triple::riscv64:
case llvm::Triple::sparc:
case llvm::Triple::sparcv9:
case llvm::Triple::sparcel:
case llvm::Triple::systemz:
case llvm::Triple::tce:
case llvm::Triple::tcele:
case llvm::Triple::thumb:
case llvm::Triple::thumbeb:
case llvm::Triple::xcore:
case llvm::Triple::le32:
case llvm::Triple::le64:
case llvm::Triple::amdil:
case llvm::Triple::amdil64:
case llvm::Triple::hsail:
case llvm::Triple::hsail64:
case llvm::Triple::spir:
case llvm::Triple::spir64:
case llvm::Triple::kalimba:
case llvm::Triple::shave:
case llvm::Triple::lanai:
case llvm::Triple::wasm32:
case llvm::Triple::wasm64:
case llvm::Triple::renderscript32:
case llvm::Triple::renderscript64:
return false;
}
}
return true;
}
bool matchesContext(CodeGenModule &CGM,
const CompleteOMPContextSelectorData &ContextData) {
for (const OMPContextSelectorData &Data : ContextData) { for (const OMPContextSelectorData &Data : ContextData) {
switch (Data.CtxSet) { switch (Data.CtxSet) {
case OMP_CTX_SET_implementation: case OMP_CTX_SET_implementation:
@ -11057,8 +11145,22 @@ bool matchesContext(const CompleteOMPContextSelectorData &ContextData) {
if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data)) if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data))
return false; return false;
break; break;
case OMP_CTX_kind:
case OMP_CTX_unknown: case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind."); llvm_unreachable(
"Unexpected context selector kind in implementation set.");
}
break;
case OMP_CTX_SET_device:
switch (Data.Ctx) {
case OMP_CTX_kind:
if (!checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(
Data, CGM))
return false;
break;
case OMP_CTX_vendor:
case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind in device set.");
} }
break; break;
case OMP_CTX_SET_unknown: case OMP_CTX_SET_unknown:
@ -11089,8 +11191,21 @@ translateAttrToContextSelectorData(ASTContext &C,
Data.back().Names = Data.back().Names =
llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end()); llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
break; break;
case OMP_CTX_kind:
case OMP_CTX_unknown: case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind."); llvm_unreachable(
"Unexpected context selector kind in implementation set.");
}
break;
case OMP_CTX_SET_device:
switch (Ctx) {
case OMP_CTX_kind:
Data.back().Names =
llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
break;
case OMP_CTX_vendor:
case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind in device set.");
} }
break; break;
case OMP_CTX_SET_unknown: case OMP_CTX_SET_unknown:
@ -11100,27 +11215,59 @@ translateAttrToContextSelectorData(ASTContext &C,
return Data; return Data;
} }
static bool isStrictSubset(const CompleteOMPContextSelectorData &LHS,
const CompleteOMPContextSelectorData &RHS) {
llvm::SmallDenseMap<std::pair<int, int>, llvm::StringSet<>, 4> RHSData;
for (const OMPContextSelectorData &D : RHS) {
auto &Pair = RHSData.FindAndConstruct(std::make_pair(D.CtxSet, D.Ctx));
Pair.getSecond().insert(D.Names.begin(), D.Names.end());
}
bool AllSetsAreEqual = true;
for (const OMPContextSelectorData &D : LHS) {
auto It = RHSData.find(std::make_pair(D.CtxSet, D.Ctx));
if (It == RHSData.end())
return false;
if (D.Names.size() > It->getSecond().size())
return false;
if (llvm::set_union(It->getSecond(), D.Names))
return false;
AllSetsAreEqual =
AllSetsAreEqual && (D.Names.size() == It->getSecond().size());
}
return LHS.size() != RHS.size() || !AllSetsAreEqual;
}
static bool greaterCtxScore(const CompleteOMPContextSelectorData &LHS, static bool greaterCtxScore(const CompleteOMPContextSelectorData &LHS,
const CompleteOMPContextSelectorData &RHS) { const CompleteOMPContextSelectorData &RHS) {
// Score is calculated as sum of all scores + 1. // Score is calculated as sum of all scores + 1.
llvm::APSInt LHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false); llvm::APSInt LHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
for (const OMPContextSelectorData &Data : LHS) { bool RHSIsSubsetOfLHS = isStrictSubset(RHS, LHS);
if (Data.Score.getBitWidth() > LHSScore.getBitWidth()) { if (RHSIsSubsetOfLHS) {
LHSScore = LHSScore.extend(Data.Score.getBitWidth()) + Data.Score; LHSScore = llvm::APSInt::get(0);
} else if (Data.Score.getBitWidth() < LHSScore.getBitWidth()) { } else {
LHSScore += Data.Score.extend(LHSScore.getBitWidth()); for (const OMPContextSelectorData &Data : LHS) {
} else { if (Data.Score.getBitWidth() > LHSScore.getBitWidth()) {
LHSScore += Data.Score; LHSScore = LHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
} else if (Data.Score.getBitWidth() < LHSScore.getBitWidth()) {
LHSScore += Data.Score.extend(LHSScore.getBitWidth());
} else {
LHSScore += Data.Score;
}
} }
} }
llvm::APSInt RHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false); llvm::APSInt RHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
for (const OMPContextSelectorData &Data : RHS) { if (!RHSIsSubsetOfLHS && isStrictSubset(LHS, RHS)) {
if (Data.Score.getBitWidth() > RHSScore.getBitWidth()) { RHSScore = llvm::APSInt::get(0);
RHSScore = RHSScore.extend(Data.Score.getBitWidth()) + Data.Score; } else {
} else if (Data.Score.getBitWidth() < RHSScore.getBitWidth()) { for (const OMPContextSelectorData &Data : RHS) {
RHSScore += Data.Score.extend(RHSScore.getBitWidth()); if (Data.Score.getBitWidth() > RHSScore.getBitWidth()) {
} else { RHSScore = RHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
RHSScore += Data.Score; } else if (Data.Score.getBitWidth() < RHSScore.getBitWidth()) {
RHSScore += Data.Score.extend(RHSScore.getBitWidth());
} else {
RHSScore += Data.Score;
}
} }
} }
return llvm::APSInt::compareValues(LHSScore, RHSScore) >= 0; return llvm::APSInt::compareValues(LHSScore, RHSScore) >= 0;
@ -11128,7 +11275,7 @@ static bool greaterCtxScore(const CompleteOMPContextSelectorData &LHS,
/// Finds the variant function that matches current context with its context /// Finds the variant function that matches current context with its context
/// selector. /// selector.
static const FunctionDecl *getDeclareVariantFunction(ASTContext &Ctx, static const FunctionDecl *getDeclareVariantFunction(CodeGenModule &CGM,
const FunctionDecl *FD) { const FunctionDecl *FD) {
if (!FD->hasAttrs() || !FD->hasAttr<OMPDeclareVariantAttr>()) if (!FD->hasAttrs() || !FD->hasAttr<OMPDeclareVariantAttr>())
return FD; return FD;
@ -11137,8 +11284,8 @@ static const FunctionDecl *getDeclareVariantFunction(ASTContext &Ctx,
CompleteOMPContextSelectorData TopMostData; CompleteOMPContextSelectorData TopMostData;
for (const auto *A : FD->specific_attrs<OMPDeclareVariantAttr>()) { for (const auto *A : FD->specific_attrs<OMPDeclareVariantAttr>()) {
CompleteOMPContextSelectorData Data = CompleteOMPContextSelectorData Data =
translateAttrToContextSelectorData(Ctx, A); translateAttrToContextSelectorData(CGM.getContext(), A);
if (!matchesContext(Data)) if (!matchesContext(CGM, Data))
continue; continue;
// If the attribute matches the context, find the attribute with the highest // If the attribute matches the context, find the attribute with the highest
// score. // score.
@ -11161,7 +11308,7 @@ bool CGOpenMPRuntime::emitDeclareVariant(GlobalDecl GD, bool IsForDefinition) {
llvm::GlobalValue *Orig = CGM.GetGlobalValue(MangledName); llvm::GlobalValue *Orig = CGM.GetGlobalValue(MangledName);
if (Orig && !Orig->isDeclaration()) if (Orig && !Orig->isDeclaration())
return false; return false;
const FunctionDecl *NewFD = getDeclareVariantFunction(CGM.getContext(), D); const FunctionDecl *NewFD = getDeclareVariantFunction(CGM, D);
// Emit original function if it does not have declare variant attribute or the // Emit original function if it does not have declare variant attribute or the
// context does not match. // context does not match.
if (NewFD == D) if (NewFD == D)

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

@ -874,6 +874,7 @@ parseImplementationSelector(Parser &P, SourceLocation Loc,
Data.emplace_back(OMP_CTX_SET_implementation, CSKind, Score, Vendors); Data.emplace_back(OMP_CTX_SET_implementation, CSKind, Score, Vendors);
break; break;
} }
case OMP_CTX_kind:
case OMP_CTX_unknown: case OMP_CTX_unknown:
P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected) P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
<< "implementation"; << "implementation";
@ -885,6 +886,91 @@ parseImplementationSelector(Parser &P, SourceLocation Loc,
} }
} }
/// Parse context selector for 'device' selector set:
/// 'kind' '(' <kind> { ',' <kind> } ')'
static void
parseDeviceSelector(Parser &P, SourceLocation Loc,
llvm::StringMap<SourceLocation> &UsedCtx,
SmallVectorImpl<Sema::OMPCtxSelectorData> &Data) {
const Token &Tok = P.getCurToken();
// Parse inner context selector set name, if any.
if (!Tok.is(tok::identifier)) {
P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
<< "device";
// Skip until either '}', ')', or end of directive.
while (!P.SkipUntil(tok::r_brace, tok::r_paren,
tok::annot_pragma_openmp_end, Parser::StopBeforeMatch))
;
return;
}
Sema::OMPCtxStringType Buffer;
StringRef CtxSelectorName = P.getPreprocessor().getSpelling(Tok, Buffer);
auto Res = UsedCtx.try_emplace(CtxSelectorName, Tok.getLocation());
if (!Res.second) {
// OpenMP 5.0, 2.3.2 Context Selectors, Restrictions.
// Each trait-selector-name can only be specified once.
P.Diag(Tok.getLocation(), diag::err_omp_declare_variant_ctx_mutiple_use)
<< CtxSelectorName << "device";
P.Diag(Res.first->getValue(), diag::note_omp_declare_variant_ctx_used_here)
<< CtxSelectorName;
}
OpenMPContextSelectorKind CSKind = getOpenMPContextSelector(CtxSelectorName);
(void)P.ConsumeToken();
switch (CSKind) {
case OMP_CTX_kind: {
// Parse '('.
BalancedDelimiterTracker T(P, tok::l_paren, tok::annot_pragma_openmp_end);
(void)T.expectAndConsume(diag::err_expected_lparen_after,
CtxSelectorName.data());
llvm::UniqueVector<Sema::OMPCtxStringType> Kinds;
do {
// Parse <kind>.
StringRef KindName;
if (Tok.is(tok::identifier)) {
Buffer.clear();
KindName = P.getPreprocessor().getSpelling(P.getCurToken(), Buffer);
SourceLocation SLoc = P.getCurToken().getLocation();
(void)P.ConsumeToken();
if (llvm::StringSwitch<bool>(KindName)
.Case("host", false)
.Case("nohost", false)
.Case("cpu", false)
.Case("gpu", false)
.Case("fpga", false)
.Default(true)) {
P.Diag(SLoc, diag::err_omp_wrong_device_kind_trait) << KindName;
} else {
Kinds.insert(KindName);
}
} else {
P.Diag(Tok.getLocation(), diag::err_omp_declare_variant_item_expected)
<< "'host', 'nohost', 'cpu', 'gpu', or 'fpga'"
<< "kind"
<< "device";
}
if (!P.TryConsumeToken(tok::comma) && Tok.isNot(tok::r_paren)) {
P.Diag(Tok, diag::err_expected_punc)
<< (KindName.empty() ? "kind of device" : KindName);
}
} while (Tok.is(tok::identifier));
// Parse ')'.
(void)T.consumeClose();
if (!Kinds.empty())
Data.emplace_back(OMP_CTX_SET_device, CSKind, ExprResult(), Kinds);
break;
}
case OMP_CTX_vendor:
case OMP_CTX_unknown:
P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
<< "device";
// Skip until either '}', ')', or end of directive.
while (!P.SkipUntil(tok::r_brace, tok::r_paren,
tok::annot_pragma_openmp_end, Parser::StopBeforeMatch))
;
return;
}
}
/// Parses clauses for 'declare variant' directive. /// Parses clauses for 'declare variant' directive.
/// clause: /// clause:
/// <selector_set_name> '=' '{' <context_selectors> '}' /// <selector_set_name> '=' '{' <context_selectors> '}'
@ -935,6 +1021,9 @@ bool Parser::parseOpenMPContextSelectors(
case OMP_CTX_SET_implementation: case OMP_CTX_SET_implementation:
parseImplementationSelector(*this, Loc, UsedCtx, Data); parseImplementationSelector(*this, Loc, UsedCtx, Data);
break; break;
case OMP_CTX_SET_device:
parseDeviceSelector(*this, Loc, UsedCtx, Data);
break;
case OMP_CTX_SET_unknown: case OMP_CTX_SET_unknown:
// Skip until either '}', ')', or end of directive. // Skip until either '}', ')', or end of directive.
while (!SkipUntil(tok::r_brace, tok::r_paren, while (!SkipUntil(tok::r_brace, tok::r_paren,

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

@ -5364,7 +5364,7 @@ void Sema::ActOnOpenMPDeclareVariantDirective(
SmallVector<Expr *, 4> CtxScores; SmallVector<Expr *, 4> CtxScores;
SmallVector<unsigned, 4> CtxSets; SmallVector<unsigned, 4> CtxSets;
SmallVector<unsigned, 4> Ctxs; SmallVector<unsigned, 4> Ctxs;
SmallVector<StringRef, 4> ImplVendors; SmallVector<StringRef, 4> ImplVendors, DeviceKinds;
bool IsError = false; bool IsError = false;
for (const OMPCtxSelectorData &D : Data) { for (const OMPCtxSelectorData &D : Data) {
OpenMPContextSelectorSetKind CtxSet = D.CtxSet; OpenMPContextSelectorSetKind CtxSet = D.CtxSet;
@ -5384,7 +5384,19 @@ void Sema::ActOnOpenMPDeclareVariantDirective(
Score = VerifyIntegerConstantExpression(Score).get(); Score = VerifyIntegerConstantExpression(Score).get();
} }
} else { } else {
Score = ActOnIntegerConstant(SourceLocation(), 0).get(); // OpenMP 5.0, 2.3.3 Matching and Scoring Context Selectors.
// The kind, arch, and isa selectors are given the values 2^l, 2^(l+1) and
// 2^(l+2), respectively, where l is the number of traits in the construct
// set.
// TODO: implement correct logic for isa and arch traits.
// TODO: take the construct context set into account when it is
// implemented.
int L = 0; // Currently set the number of traits in construct set to 0,
// since the construct trait set in not supported yet.
if (CtxSet == OMP_CTX_SET_device && Ctx == OMP_CTX_kind)
Score = ActOnIntegerConstant(SourceLocation(), std::pow(2, L)).get();
else
Score = ActOnIntegerConstant(SourceLocation(), 0).get();
} }
switch (CtxSet) { switch (CtxSet) {
case OMP_CTX_SET_implementation: case OMP_CTX_SET_implementation:
@ -5392,6 +5404,17 @@ void Sema::ActOnOpenMPDeclareVariantDirective(
case OMP_CTX_vendor: case OMP_CTX_vendor:
ImplVendors.append(D.Names.begin(), D.Names.end()); ImplVendors.append(D.Names.begin(), D.Names.end());
break; break;
case OMP_CTX_kind:
case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind.");
}
break;
case OMP_CTX_SET_device:
switch (Ctx) {
case OMP_CTX_kind:
DeviceKinds.append(D.Names.begin(), D.Names.end());
break;
case OMP_CTX_vendor:
case OMP_CTX_unknown: case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind."); llvm_unreachable("Unexpected context selector kind.");
} }
@ -5408,7 +5431,8 @@ void Sema::ActOnOpenMPDeclareVariantDirective(
auto *NewAttr = OMPDeclareVariantAttr::CreateImplicit( auto *NewAttr = OMPDeclareVariantAttr::CreateImplicit(
Context, VariantRef, CtxScores.begin(), CtxScores.size(), Context, VariantRef, CtxScores.begin(), CtxScores.size(),
CtxSets.begin(), CtxSets.size(), Ctxs.begin(), Ctxs.size(), CtxSets.begin(), CtxSets.size(), Ctxs.begin(), Ctxs.size(),
ImplVendors.begin(), ImplVendors.size(), SR); ImplVendors.begin(), ImplVendors.size(), DeviceKinds.begin(),
DeviceKinds.size(), SR);
FD->addAttr(NewAttr); FD->addAttr(NewAttr);
} }
} }

11
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
View File

@ -410,6 +410,17 @@ static void instantiateOMPDeclareVariantAttr(
case OMP_CTX_vendor: case OMP_CTX_vendor:
Data.emplace_back(CtxSet, Ctx, Score, Attr.implVendors()); Data.emplace_back(CtxSet, Ctx, Score, Attr.implVendors());
break; break;
case OMP_CTX_kind:
case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind.");
}
break;
case OMP_CTX_SET_device:
switch (Ctx) {
case OMP_CTX_kind:
Data.emplace_back(CtxSet, Ctx, Score, Attr.deviceKinds());
break;
case OMP_CTX_vendor:
case OMP_CTX_unknown: case OMP_CTX_unknown:
llvm_unreachable("Unexpected context selector kind."); llvm_unreachable("Unexpected context selector kind.");
} }

16
clang/test/OpenMP/declare_variant_ast_print.c
View File

@ -8,15 +8,19 @@ int foo(void);
#pragma omp declare variant(foo) match(xxx={}, yyy={ccc}) #pragma omp declare variant(foo) match(xxx={}, yyy={ccc})
#pragma omp declare variant(foo) match(xxx={vvv}) #pragma omp declare variant(foo) match(xxx={vvv})
#pragma omp declare variant(foo) match(implementation={vendor(llvm)}) #pragma omp declare variant(foo) match(implementation={vendor(llvm)}, device={kind(fpga)})
#pragma omp declare variant(foo) match(implementation={vendor(llvm), xxx}) #pragma omp declare variant(foo) match(implementation={vendor(llvm), xxx})
#pragma omp declare variant(foo) match(implementation={vendor(unknown)}) #pragma omp declare variant(foo) match(implementation={vendor(unknown)}, device={kind(gpu)})
#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)}) #pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)}, device={kind(cpu, nohost)})
#pragma omp declare variant(foo) match(device={kind(host)})
#pragma omp declare variant(foo) match(device={kind(nohost), xxx})
int bar(void); int bar(void);
// CHECK: int foo(); // CHECK: int foo();
// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(5):ibm, xxx)}) // CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(nohost)})
// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):unknown)}) // CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(host)})
// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(5):ibm, xxx)},device={kind(cpu, nohost)})
// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):unknown)},device={kind(gpu)})
// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)})
// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
// CHECK-NEXT: int bar(); // CHECK-NEXT: int bar();

80
clang/test/OpenMP/declare_variant_ast_print.cpp
View File

@ -17,20 +17,20 @@ T foofoo() { return T(); }
// CHECK-NEXT: return int(); // CHECK-NEXT: return int();
// CHECK-NEXT: } // CHECK-NEXT: }
// CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5):ibm)}) // CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5):ibm)},device={kind(fpga)})
// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)}) // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)})
// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
// CHECK-NEXT: int bar(); // CHECK-NEXT: int bar();
#pragma omp declare variant(foofoo <int>) match(xxx = {}) #pragma omp declare variant(foofoo <int>) match(xxx = {})
#pragma omp declare variant(foofoo <int>) match(xxx = {vvv}) #pragma omp declare variant(foofoo <int>) match(xxx = {vvv})
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(llvm), xxx}) #pragma omp declare variant(foofoo <int>) match(implementation={vendor(llvm), xxx}, device={kind(cpu)})
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(unknown)}) #pragma omp declare variant(foofoo <int>) match(implementation={vendor(unknown)})
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)}) #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga)})
int bar(); int bar();
// CHECK: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)}) // CHECK: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)},device={kind(cpu, host)})
// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):unknown)}) // CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):unknown)})
// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
// CHECK-NEXT: template <typename T, int C> T barbar(); // CHECK-NEXT: template <typename T, int C> T barbar();
#pragma omp declare variant(foofoo <T>) match(xxx = {}) #pragma omp declare variant(foofoo <T>) match(xxx = {})
#pragma omp declare variant(foofoo <T>) match(xxx = {vvv}) #pragma omp declare variant(foofoo <T>) match(xxx = {vvv})
@ -38,15 +38,15 @@ int bar();
#pragma omp declare variant(foofoo <T>) match(user = {score(<expr>) : condition(<expr>)}) #pragma omp declare variant(foofoo <T>) match(user = {score(<expr>) : condition(<expr>)})
#pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)}) #pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)})
#pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)}) #pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)})
#pragma omp declare variant(foofoo <T>) match(implementation={vendor(llvm)}) #pragma omp declare variant(foofoo <T>) match(implementation={vendor(llvm)},device={kind(cpu)})
#pragma omp declare variant(foofoo <T>) match(implementation={vendor(unknown)}) #pragma omp declare variant(foofoo <T>) match(implementation={vendor(unknown)})
#pragma omp declare variant(foofoo <T>) match(implementation={vendor(score(C+5): ibm, xxx, ibm)}) #pragma omp declare variant(foofoo <T>) match(implementation={vendor(score(C+5): ibm, xxx, ibm)},device={kind(cpu,host)})
template <typename T, int C> template <typename T, int C>
T barbar(); T barbar();
// CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)}) // CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)},device={kind(cpu, host)})
// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)}) // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)})
// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
// CHECK-NEXT: template<> int barbar<int, 3>(); // CHECK-NEXT: template<> int barbar<int, 3>();
// CHECK-NEXT: int baz() { // CHECK-NEXT: int baz() {
@ -66,19 +66,19 @@ template <class C>
void h_ref(C *hp, C *hp2, C *hq, C *lin) { void h_ref(C *hp, C *hp2, C *hq, C *lin) {
} }
// CHECK: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):unknown)}) // CHECK: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
// CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):llvm)},device={kind(gpu)})
// CHECK-NEXT: template <class C> void h(C *hp, C *hp2, C *hq, C *lin) { // CHECK-NEXT: template <class C> void h(C *hp, C *hp2, C *hq, C *lin) {
// CHECK-NEXT: } // CHECK-NEXT: }
#pragma omp declare variant(h_ref <C>) match(xxx = {}) #pragma omp declare variant(h_ref <C>) match(xxx = {})
#pragma omp declare variant(h_ref <C>) match(implementation={vendor(llvm)}) #pragma omp declare variant(h_ref <C>) match(implementation={vendor(llvm)}, device={kind(gpu)})
#pragma omp declare variant(h_ref <C>) match(implementation={vendor(unknown)}) #pragma omp declare variant(h_ref <C>) match(implementation={vendor(unknown)},device={kind(nohost)})
template <class C> template <class C>
void h(C *hp, C *hp2, C *hq, C *lin) { void h(C *hp, C *hp2, C *hq, C *lin) {
} }
// CHECK: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):unknown)}) // CHECK: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
// CHECK-NEXT: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):llvm)},device={kind(gpu)})
// CHECK-NEXT: template<> void h<float>(float *hp, float *hp2, float *hq, float *lin) { // CHECK-NEXT: template<> void h<float>(float *hp, float *hp2, float *hq, float *lin) {
// CHECK-NEXT: } // CHECK-NEXT: }
@ -86,7 +86,7 @@ void h(C *hp, C *hp2, C *hq, C *lin) {
// CHECK-NEXT: h((float *)hp, (float *)hp2, (float *)hq, (float *)lin); // CHECK-NEXT: h((float *)hp, (float *)hp2, (float *)hq, (float *)lin);
// CHECK-NEXT: } // CHECK-NEXT: }
#pragma omp declare variant(h_ref <double>) match(xxx = {}) #pragma omp declare variant(h_ref <double>) match(xxx = {})
#pragma omp declare variant(h_ref <double>) match(implementation={vendor(ibm)}) #pragma omp declare variant(h_ref <double>) match(implementation={vendor(ibm)},device={kind(cpu,gpu)})
#pragma omp declare variant(h_ref <double>) match(implementation={vendor(unknown)}) #pragma omp declare variant(h_ref <double>) match(implementation={vendor(unknown)})
template <> template <>
void h(double *hp, double *hp2, double *hq, double *lin) { void h(double *hp, double *hp2, double *hq, double *lin) {
@ -97,36 +97,36 @@ void h(double *hp, double *hp2, double *hq, double *lin) {
int fn(); int fn();
// CHECK: int fn(int); // CHECK: int fn(int);
int fn(int); int fn(int);
// CHECK: #pragma omp declare variant(fn) match(implementation={vendor(score(0):unknown)}) // CHECK: #pragma omp declare variant(fn) match(implementation={vendor(score(0):unknown)},device={kind(cpu, gpu)})
// CHECK-NEXT: #pragma omp declare variant(fn) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(fn) match(implementation={vendor(score(0):llvm)})
// CHECK-NEXT: int overload(); // CHECK-NEXT: int overload();
#pragma omp declare variant(fn) match(xxx = {}) #pragma omp declare variant(fn) match(xxx = {})
#pragma omp declare variant(fn) match(implementation={vendor(llvm)}) #pragma omp declare variant(fn) match(implementation={vendor(llvm)})
#pragma omp declare variant(fn) match(implementation={vendor(unknown)}) #pragma omp declare variant(fn) match(implementation={vendor(unknown)},device={kind(cpu,gpu)})
int overload(void); int overload(void);
// CHECK: int fn_deduced_variant() { // CHECK: int fn_deduced_variant() {
// CHECK-NEXT: return 0; // CHECK-NEXT: return 0;
// CHECK-NEXT: } // CHECK-NEXT: }
auto fn_deduced_variant() { return 0; } auto fn_deduced_variant() { return 0; }
// CHECK: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):unknown)}) // CHECK: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):unknown)},device={kind(gpu, nohost)})
// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):llvm)},device={kind(cpu, host)})
// CHECK-NEXT: int fn_deduced(); // CHECK-NEXT: int fn_deduced();
#pragma omp declare variant(fn_deduced_variant) match(xxx = {}) #pragma omp declare variant(fn_deduced_variant) match(xxx = {})
#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(llvm)}) #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(llvm)},device={kind(cpu,host)})
#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(unknown)}) #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(unknown)},device={kind(gpu,nohost)})
int fn_deduced(); int fn_deduced();
// CHECK: int fn_deduced_variant1(); // CHECK: int fn_deduced_variant1();
int fn_deduced_variant1(); int fn_deduced_variant1();
// CHECK: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):unknown)}) // CHECK: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):ibm)}) // CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):ibm)},device={kind(gpu, nohost)})
// CHECK-NEXT: int fn_deduced1() { // CHECK-NEXT: int fn_deduced1() {
// CHECK-NEXT: return 0; // CHECK-NEXT: return 0;
// CHECK-NEXT: } // CHECK-NEXT: }
#pragma omp declare variant(fn_deduced_variant1) match(xxx = {}) #pragma omp declare variant(fn_deduced_variant1) match(xxx = {})
#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(ibm)}) #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(ibm)},device={kind(gpu,nohost)})
#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(unknown)}) #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(unknown)},device={kind(cpu,host)})
auto fn_deduced1() { return 0; } auto fn_deduced1() { return 0; }
// CHECK: struct SpecialFuncs { // CHECK: struct SpecialFuncs {
@ -140,11 +140,11 @@ auto fn_deduced1() { return 0; }
// CHECK-NEXT: } // CHECK-NEXT: }
// CHECK-NEXT: void bar(int) { // CHECK-NEXT: void bar(int) {
// CHECK-NEXT: } // CHECK-NEXT: }
// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)}) // CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(score(0):ibm)}) // CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(score(0):ibm)},device={kind(cpu)})
// CHECK-NEXT: void foo1() { // CHECK-NEXT: void foo1() {
// CHECK-NEXT: } // CHECK-NEXT: }
// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)}) // CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
// CHECK-NEXT: void xxx(); // CHECK-NEXT: void xxx();
// CHECK-NEXT: } s; // CHECK-NEXT: } s;
struct SpecialFuncs { struct SpecialFuncs {
@ -157,14 +157,14 @@ struct SpecialFuncs {
void bar(int) {} void bar(int) {}
#pragma omp declare variant(SpecialFuncs::baz) match(xxx = {}) #pragma omp declare variant(SpecialFuncs::baz) match(xxx = {})
#pragma omp declare variant(SpecialFuncs::bar) match(xxx = {}) #pragma omp declare variant(SpecialFuncs::bar) match(xxx = {})
#pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(ibm)}) #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(ibm)},device={kind(cpu)})
#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)}) #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(nohost)})
void foo1() {} void foo1() {}
#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)}) #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(cpu, host)})
void xxx(); void xxx();
} s; } s;
// CHECK: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)}) // CHECK: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
// CHECK-NEXT: void SpecialFuncs::xxx() { // CHECK-NEXT: void SpecialFuncs::xxx() {
// CHECK-NEXT: } // CHECK-NEXT: }
void SpecialFuncs::xxx() {} void SpecialFuncs::xxx() {}
@ -173,11 +173,11 @@ void SpecialFuncs::xxx() {}
// CHECK-NEXT: } // CHECK-NEXT: }
static void static_f_variant() {} static void static_f_variant() {}
// CHECK: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):unknown)}) // CHECK: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):unknown)})
// CHECK-NEXT: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):llvm)}) // CHECK-NEXT: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
// CHECK-NEXT: static void static_f() { // CHECK-NEXT: static void static_f() {
// CHECK-NEXT: } // CHECK-NEXT: }
#pragma omp declare variant(static_f_variant) match(xxx = {}) #pragma omp declare variant(static_f_variant) match(xxx = {})
#pragma omp declare variant(static_f_variant) match(implementation={vendor(llvm)}) #pragma omp declare variant(static_f_variant) match(implementation={vendor(llvm)},device={kind(fpga)})
#pragma omp declare variant(static_f_variant) match(implementation={vendor(unknown)}) #pragma omp declare variant(static_f_variant) match(implementation={vendor(unknown)})
static void static_f() {} static void static_f() {}
@ -192,19 +192,19 @@ void bazzzz() {
// CHECK: int fn_linkage_variant(); // CHECK: int fn_linkage_variant();
// CHECK: extern "C" { // CHECK: extern "C" {
// CHECK: #pragma omp declare variant(fn_linkage_variant) match(implementation={vendor(score(0):xxx)}) // CHECK: #pragma omp declare variant(fn_linkage_variant) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
// CHECK: int fn_linkage(); // CHECK: int fn_linkage();
// CHECK: } // CHECK: }
int fn_linkage_variant(); int fn_linkage_variant();
extern "C" { extern "C" {
#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)}) #pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
int fn_linkage(); int fn_linkage();
} }
// CHECK: extern "C" int fn_linkage_variant1() // CHECK: extern "C" int fn_linkage_variant1()
// CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)}) // CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
// CHECK: int fn_linkage1(); // CHECK: int fn_linkage1();
extern "C" int fn_linkage_variant1(); extern "C" int fn_linkage_variant1();
#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)}) #pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
int fn_linkage1(); int fn_linkage1();

187
clang/test/OpenMP/declare_variant_device_kind_codegen.cpp Normal file
View File

@ -0,0 +1,187 @@
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DHOST | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DHOST
// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DHOST | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DHOST | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DHOST
// RUN: %clang_cc1 -fopenmp -x c++ -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DHOST | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DHOST | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DHOST
// RUN: %clang_cc1 -fopenmp -x c++ -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DHOST | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCPU | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCPU
// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCPU | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCPU | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCPU
// RUN: %clang_cc1 -fopenmp -x c++ -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCPU | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCPU | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCPU
// RUN: %clang_cc1 -fopenmp -x c++ -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCPU | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=x86_64-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DCPU
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DCPU | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DCPU
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DCPU | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -fopenmp-targets=ppc64le-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DCPU
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DCPU | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DCPU
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DCPU | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=x86_64-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DNOHOST
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DNOHOST | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DNOHOST
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DNOHOST | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -fopenmp-targets=ppc64le-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DNOHOST
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DNOHOST | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DNOHOST
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DNOHOST | FileCheck %s
// expected-no-diagnostics
// CHECK-NOT: ret i32 {{1|4|81|84}}
// CHECK-DAG: @_Z3barv = {{.*}}alias i32 (), i32 ()* @_Z3foov
// CHECK-DAG: @_ZN16SpecSpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
// CHECK-DAG: @_ZN16SpecSpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
// CHECK-DAG: @_ZN12SpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
// CHECK-DAG: @_Z5prio_v = {{.*}}alias i32 (), i32 ()* @_Z5prio1v
// CHECK-DAG: @_ZL6prio1_v = internal alias i32 (), i32 ()* @_ZL5prio2v
// CHECK-DAG: @_Z4callv = {{.*}}alias i32 (), i32 ()* @_Z4testv
// CHECK-DAG: @_ZL9stat_usedv = internal alias i32 (), i32 ()* @_ZL10stat_used_v
// CHECK-DAG: @_ZN12SpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
// CHECK-DAG: @fn_linkage = {{.*}}alias i32 (), i32 ()* @_Z18fn_linkage_variantv
// CHECK-DAG: @_Z11fn_linkage1v = {{.*}}alias i32 (), i32 ()* @fn_linkage_variant1
// CHECK-DAG: declare {{.*}}i32 @_Z5bazzzv()
// CHECK-DAG: declare {{.*}}i32 @_Z3bazv()
// CHECK-DAG: ret i32 2
// CHECK-DAG: ret i32 3
// CHECK-DAG: ret i32 5
// CHECK-DAG: ret i32 6
// CHECK-DAG: ret i32 7
// CHECK-DAG: ret i32 82
// CHECK-DAG: ret i32 83
// CHECK-DAG: ret i32 85
// CHECK-DAG: ret i32 86
// CHECK-DAG: ret i32 87
// CHECK-NOT: ret i32 {{1|4|81|84}}
#ifndef HEADER
#define HEADER
#pragma omp declare target
#ifdef HOST
#define CORRECT host
#define SUBSET host, cpu
#define WRONG host, nohost
#endif // HOST
#ifdef CPU
#define CORRECT cpu
#define SUBSET host, cpu
#define WRONG cpu, gpu
#endif // CPU
#ifdef NOHOST
#define CORRECT nohost
#define SUBSET nohost, cpu
#define WRONG nohost, host
#endif // NOHOST
int foo() { return 2; }
#pragma omp declare variant(foo) match(device = {kind(CORRECT)})
int bar() { return 1; }
int bazzz();
#pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})
int baz() { return 1; }
int test();
#pragma omp declare variant(test) match(device = {kind(CORRECT)})
int call() { return 1; }
static int stat_unused_();
#pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})
static int stat_unused() { return 1; }
static int stat_used_();
#pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)})
static int stat_used() { return 1; }
int main() { return bar() + baz() + call() + stat_used(); }
int test() { return 3; }
static int stat_unused_() { return 4; }
static int stat_used_() { return 5; }
struct SpecialFuncs {
void vd() {}
SpecialFuncs();
~SpecialFuncs();
int method_() { return 6; }
#pragma omp declare variant(SpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int method() { return 1; }
#pragma omp declare variant(SpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int Method();
} s;
int SpecialFuncs::Method() { return 1; }
struct SpecSpecialFuncs {
void vd() {}
SpecSpecialFuncs();
~SpecSpecialFuncs();
int method_();
#pragma omp declare variant(SpecSpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int method() { return 1; }
#pragma omp declare variant(SpecSpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int Method();
} s1;
int SpecSpecialFuncs::method_() { return 7; }
int SpecSpecialFuncs::Method() { return 1; }
void xxx() {
(void)s.method();
(void)s1.method();
}
int prio() { return 81; }
int prio1() { return 82; }
#pragma omp declare variant(prio) match(device = {kind(SUBSET)})
#pragma omp declare variant(prio1) match(device = {kind(CORRECT)})
int prio_() { return 1; }
static int prio2() { return 83; }
static int prio3() { return 84; }
static int prio4() { return 84; }
#pragma omp declare variant(prio4) match(device = {kind(SUBSET)})
#pragma omp declare variant(prio2) match(device = {kind(CORRECT)})
#pragma omp declare variant(prio3) match(device = {kind(SUBSET)})
static int prio1_() { return 1; }
int int_fn() { return prio1_(); }
int fn_linkage_variant() { return 85; }
extern "C" {
#pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)})
int fn_linkage() { return 1; }
}
extern "C" int fn_linkage_variant1() { return 86; }
#pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})
int fn_linkage1() { return 1; }
int fn_variant2() { return 1; }
#pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})
int fn2() { return 87; }
#pragma omp end declare target
#endif // HEADER

13
clang/test/OpenMP/declare_variant_messages.c
View File

@ -35,6 +35,17 @@ int foo(void);
#pragma omp declare variant(foo) match(implementation={vendor(score(2 ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}} #pragma omp declare variant(foo) match(implementation={vendor(score(2 ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}}
#pragma omp declare variant(foo) match(implementation={vendor(score(foo()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integer constant expression}} #pragma omp declare variant(foo) match(implementation={vendor(score(foo()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integer constant expression}}
#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}} #pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}}
#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm), kind(cpu)}) // expected-warning {{unknown context selector in 'implementation' context selector set of 'omp declare variant' directive, ignored}}
#pragma omp declare variant(foo) match(device={xxx}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}}
#pragma omp declare variant(foo) match(device={kind}) // expected-error {{expected '(' after 'kind'}} expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
#pragma omp declare variant(foo) match(device={kind(}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error 2 {{expected ')'}} expected-note {{to match this '('}}
#pragma omp declare variant(foo) match(device={kind()}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}}
#pragma omp declare variant(foo) match(device={kind(score cpu)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foo) match(device={kind(score( ibm)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foo) match(device={kind(score(2 gpu)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foo) match(device={kind(score(foo()) ibm)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foo) match(device={kind(score(5): host), kind(llvm)}) // expected-error {{context trait selector 'kind' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'kind' used here}} expected-error {{expected ')' or ',' after 'score'}} expected-note {{to match this '('}} expected-error {{expected ')'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}} expected-error {{unknown 'llvm' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foo) match(device={kind(score(5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}}
int bar(void); int bar(void);
// expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}} // expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}}
@ -105,7 +116,7 @@ void diff_ret(void);
void marked(void); void marked(void);
void not_marked(void); void not_marked(void);
// expected-note@+1 {{marked as 'declare variant' here}} // expected-note@+1 {{marked as 'declare variant' here}}
#pragma omp declare variant(not_marked) match(implementation={vendor(unknown)}) #pragma omp declare variant(not_marked) match(implementation={vendor(unknown)}, device={kind(cpu)})
void marked_variant(void); void marked_variant(void);
// expected-warning@+1 {{variant function in '#pragma omp declare variant' is itself marked as '#pragma omp declare variant'}} // expected-warning@+1 {{variant function in '#pragma omp declare variant' is itself marked as '#pragma omp declare variant'}}
#pragma omp declare variant(marked_variant) match(xxx={}) #pragma omp declare variant(marked_variant) match(xxx={})

22
clang/test/OpenMP/declare_variant_messages.cpp
View File

@ -38,6 +38,17 @@ T foofoo(); // expected-note 2 {{declared here}}
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(2 ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}} #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(2 ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}}
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(foofoo <int>()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integral constant expression}} expected-note {{non-constexpr function 'foofoo<int>' cannot be used in a constant expression}} #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(foofoo <int>()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integral constant expression}} expected-note {{non-constexpr function 'foofoo<int>' cannot be used in a constant expression}}
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}} #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}}
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm), kind(cpu)}) // expected-warning {{unknown context selector in 'implementation' context selector set of 'omp declare variant' directive, ignored}}
#pragma omp declare variant(foofoo <int>) match(device={xxx}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}}
#pragma omp declare variant(foofoo <int>) match(device={kind}) // expected-error {{expected '(' after 'kind'}} expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
#pragma omp declare variant(foofoo <int>) match(device={kind(}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error 2 {{expected ')'}} expected-note {{to match this '('}}
#pragma omp declare variant(foofoo <int>) match(device={kind()}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score cpu)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score( ibm)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(2 gpu)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(foofoo <int>()) ibm)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(5): host), kind(llvm)}) // expected-error {{context trait selector 'kind' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'kind' used here}} expected-error {{expected ')' or ',' after 'score'}} expected-note {{to match this '('}} expected-error {{expected ')'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}} expected-error {{unknown 'llvm' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
int bar(); int bar();
#pragma omp declare variant // expected-error {{expected '(' after 'declare variant'}} #pragma omp declare variant // expected-error {{expected '(' after 'declare variant'}}
@ -67,6 +78,17 @@ int bar();
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(C ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}} #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(C ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}}
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(foofoo <int>()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integral constant expression}} expected-note {{non-constexpr function 'foofoo<int>' cannot be used in a constant expression}} #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(foofoo <int>()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integral constant expression}} expected-note {{non-constexpr function 'foofoo<int>' cannot be used in a constant expression}}
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(C+5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}} #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(C+5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}}
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm), kind(cpu)}) // expected-warning {{unknown context selector in 'implementation' context selector set of 'omp declare variant' directive, ignored}}
#pragma omp declare variant(foofoo <int>) match(device={xxx}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}}
#pragma omp declare variant(foofoo <int>) match(device={kind}) // expected-error {{expected '(' after 'kind'}} expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
#pragma omp declare variant(foofoo <int>) match(device={kind(}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error 2 {{expected ')'}} expected-note {{to match this '('}}
#pragma omp declare variant(foofoo <int>) match(device={kind()}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score cpu)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score( ibm)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(C gpu)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(foofoo <int>()) ibm)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(C+5): host), kind(llvm)}) // expected-error {{context trait selector 'kind' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'kind' used here}} expected-error {{expected ')' or ',' after 'score'}} expected-note {{to match this '('}} expected-error {{expected ')'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}} expected-error {{unknown 'llvm' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(C+5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
template <typename T, int C> template <typename T, int C>
T barbar(); T barbar();

137
clang/test/OpenMP/declare_variant_mixed_codegen.cpp Normal file
View File

@ -0,0 +1,137 @@
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s
// expected-no-diagnostics
// CHECK-NOT: ret i32 {{1|4|81|84}}
// CHECK-DAG: @_Z3barv = {{.*}}alias i32 (), i32 ()* @_Z3foov
// CHECK-DAG: @_ZN16SpecSpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
// CHECK-DAG: @_ZN16SpecSpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
// CHECK-DAG: @_ZN12SpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
// CHECK-DAG: @_Z5prio_v = {{.*}}alias i32 (), i32 ()* @_Z5prio1v
// CHECK-DAG: @_ZL6prio1_v = internal alias i32 (), i32 ()* @_ZL5prio2v
// CHECK-DAG: @_Z4callv = {{.*}}alias i32 (), i32 ()* @_Z4testv
// CHECK-DAG: @_ZL9stat_usedv = internal alias i32 (), i32 ()* @_ZL10stat_used_v
// CHECK-DAG: @_ZN12SpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
// CHECK-DAG: @fn_linkage = {{.*}}alias i32 (), i32 ()* @_Z18fn_linkage_variantv
// CHECK-DAG: @_Z11fn_linkage1v = {{.*}}alias i32 (), i32 ()* @fn_linkage_variant1
// CHECK-DAG: declare {{.*}}i32 @_Z5bazzzv()
// CHECK-DAG: declare {{.*}}i32 @_Z3bazv()
// CHECK-DAG: ret i32 2
// CHECK-DAG: ret i32 3
// CHECK-DAG: ret i32 5
// CHECK-DAG: ret i32 6
// CHECK-DAG: ret i32 7
// CHECK-DAG: ret i32 82
// CHECK-DAG: ret i32 83
// CHECK-DAG: ret i32 85
// CHECK-DAG: ret i32 86
// CHECK-DAG: ret i32 87
// CHECK-NOT: ret i32 {{1|4|81|84}}
#ifndef HEADER
#define HEADER
int foo() { return 2; }
#pragma omp declare variant(foo) match(implementation = {vendor(llvm)}, device={kind(cpu)})
int bar() { return 1; }
int bazzz();
#pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)}, device={kind(host)})
int baz() { return 1; }
int test();
#pragma omp declare variant(test) match(implementation = {vendor(llvm)}, device={kind(cpu)})
int call() { return 1; }
static int stat_unused_();
#pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)}, device={kind(cpu)})
static int stat_unused() { return 1; }
static int stat_used_();
#pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)}, device={kind(host)})
static int stat_used() { return 1; }
int main() { return bar() + baz() + call() + stat_used(); }
int test() { return 3; }
static int stat_unused_() { return 4; }
static int stat_used_() { return 5; }
struct SpecialFuncs {
void vd() {}
SpecialFuncs();
~SpecialFuncs();
int method_() { return 6; }
#pragma omp declare variant(SpecialFuncs::method_) \
match(implementation = {vendor(llvm)}, device={kind(cpu)})
int method() { return 1; }
#pragma omp declare variant(SpecialFuncs::method_) \
match(implementation = {vendor(llvm)}, device={kind(host)})
int Method();
} s;
int SpecialFuncs::Method() { return 1; }
struct SpecSpecialFuncs {
void vd() {}
SpecSpecialFuncs();
~SpecSpecialFuncs();
int method_();
#pragma omp declare variant(SpecSpecialFuncs::method_) \
match(implementation = {vendor(llvm)}, device={kind(cpu)})
int method() { return 1; }
#pragma omp declare variant(SpecSpecialFuncs::method_) \
match(implementation = {vendor(llvm)}, device={kind(host)})
int Method();
} s1;
int SpecSpecialFuncs::method_() { return 7; }
int SpecSpecialFuncs::Method() { return 1; }
void xxx() {
(void)s.method();
(void)s1.method();
}
int prio() { return 81; }
int prio1() { return 82; }
#pragma omp declare variant(prio) match(implementation = {vendor(score(2): llvm)}, device={kind(cpu,host)})
#pragma omp declare variant(prio1) match(implementation = {vendor(score(1): llvm)}, device={kind(cpu)})
int prio_() { return 1; }
static int prio2() { return 83; }
static int prio3() { return 84; }
static int prio4() { return 84; }
#pragma omp declare variant(prio4) match(implementation = {vendor(score(8): llvm)},device={kind(cpu,host)})
#pragma omp declare variant(prio2) match(implementation = {vendor(score(5): llvm)})
#pragma omp declare variant(prio3) match(implementation = {vendor(score(7): llvm)}, device={kind(cpu)})
static int prio1_() { return 1; }
int int_fn() { return prio1_(); }
int fn_linkage_variant() { return 85; }
extern "C" {
#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)}, device={kind(cpu)})
int fn_linkage() { return 1; }
}
extern "C" int fn_linkage_variant1() { return 86; }
#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)}, device={kind(host)})
int fn_linkage1() { return 1; }
int fn_variant2() { return 1; }
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm, ibm)}, device={kind(cpu)})
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(cpu,gpu)})
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(nohost)})
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(cpu,nohost)})
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(gpu)})
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(fpga)})
int fn2() { return 87; }
#endif // HEADER

170
clang/test/OpenMP/nvptx_declare_variant_device_kind_codegen.cpp Normal file
View File

@ -0,0 +1,170 @@
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DGPU
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DGPU | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -fopenmp-version=50 -DGPU
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -fopenmp-version=50 -DGPU | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DNOHOST
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -fopenmp-version=50 -DNOHOST
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -fopenmp-version=50 -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
// expected-no-diagnostics
// CHECK-NOT: ret i32 {{1|81|84}}
// CHECK-DAG: define {{.*}}i32 @_Z3barv()
// CHECK-DAG: define {{.*}}i32 @_ZN16SpecSpecialFuncs6MethodEv(%struct.SpecSpecialFuncs* %{{.+}})
// CHECK-DAG: define {{.*}}i32 @_ZN12SpecialFuncs6MethodEv(%struct.SpecialFuncs* %{{.+}})
// CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN16SpecSpecialFuncs6methodEv(%struct.SpecSpecialFuncs* %{{.+}})
// CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN12SpecialFuncs6methodEv(%struct.SpecialFuncs* %{{.+}})
// CHECK-DAG: define {{.*}}i32 @_Z5prio_v()
// CHECK-DAG: define internal i32 @_ZL6prio1_v()
// CHECK-DAG: define {{.*}}i32 @_Z4callv()
// CHECK-DAG: define internal i32 @_ZL9stat_usedv()
// CHECK-DAG: define {{.*}}i32 @fn_linkage()
// CHECK-DAG: define {{.*}}i32 @_Z11fn_linkage1v()
// CHECK-DAG: ret i32 2
// CHECK-DAG: ret i32 3
// CHECK-DAG: ret i32 4
// CHECK-DAG: ret i32 5
// CHECK-DAG: ret i32 6
// CHECK-DAG: ret i32 7
// CHECK-DAG: ret i32 82
// CHECK-DAG: ret i32 83
// CHECK-DAG: ret i32 85
// CHECK-DAG: ret i32 86
// CHECK-DAG: ret i32 87
// Outputs for function members
// CHECK-DAG: ret i32 6
// CHECK-DAG: ret i32 7
// CHECK-NOT: ret i32 {{1|81|84}}
#ifndef HEADER
#define HEADER
#ifdef GPU
#define CORRECT gpu
#define SUBSET nohost, gpu
#define WRONG cpu, gpu
#endif // GPU
#ifdef NOHOST
#define CORRECT nohost
#define SUBSET nohost, gpu
#define WRONG nohost, host
#endif // NOHOST
int foo() { return 2; }
int bazzz();
int test();
static int stat_unused_();
static int stat_used_();
#pragma omp declare target
#pragma omp declare variant(foo) match(device = {kind(CORRECT)})
int bar() { return 1; }
#pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})
int baz() { return 1; }
#pragma omp declare variant(test) match(device = {kind(CORRECT)})
int call() { return 1; }
#pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})
static int stat_unused() { return 1; }
#pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)})
static int stat_used() { return 1; }
#pragma omp end declare target
int main() {
int res;
#pragma omp target map(from \
: res)
res = bar() + baz() + call();
return res;
}
int test() { return 3; }
static int stat_unused_() { return 4; }
static int stat_used_() { return 5; }
#pragma omp declare target
struct SpecialFuncs {
void vd() {}
SpecialFuncs();
~SpecialFuncs();
int method_() { return 6; }
#pragma omp declare variant(SpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int method() { return 1; }
#pragma omp declare variant(SpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int Method();
} s;
int SpecialFuncs::Method() { return 1; }
struct SpecSpecialFuncs {
void vd() {}
SpecSpecialFuncs();
~SpecSpecialFuncs();
int method_();
#pragma omp declare variant(SpecSpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int method() { return 1; }
#pragma omp declare variant(SpecSpecialFuncs::method_) \
match(device = {kind(CORRECT)})
int Method();
} s1;
#pragma omp end declare target
int SpecSpecialFuncs::method_() { return 7; }
int SpecSpecialFuncs::Method() { return 1; }
int prio() { return 81; }
int prio1() { return 82; }
static int prio2() { return 83; }
static int prio3() { return 84; }
static int prio4() { return 84; }
int fn_linkage_variant() { return 85; }
extern "C" int fn_linkage_variant1() { return 86; }
int fn_variant2() { return 1; }
#pragma omp declare target
void xxx() {
(void)s.method();
(void)s1.method();
}
#pragma omp declare variant(prio) match(device = {kind(SUBSET)})
#pragma omp declare variant(prio1) match(device = {kind(CORRECT)})
int prio_() { return 1; }
#pragma omp declare variant(prio4) match(device = {kind(SUBSET)})
#pragma omp declare variant(prio2) match(device = {kind(CORRECT)})
#pragma omp declare variant(prio3) match(device = {kind(SUBSET)})
static int prio1_() { return 1; }
int int_fn() { return prio1_(); }
extern "C" {
#pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)})
int fn_linkage() { return 1; }
}
#pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})
int fn_linkage1() { return 1; }
#pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})
int fn2() { return 87; }
#pragma omp end declare target
#endif // HEADER