[clang] c4a9a37 - [OpenACC] device_type clause Sema for Compute constructs
via cfe-commits
cfe-commits at lists.llvm.org
Mon May 13 07:50:24 PDT 2024
Author: erichkeane
Date: 2024-05-13T07:50:19-07:00
New Revision: c4a9a374749deb5f2a932a7d4ef9321be1b2ae5d
URL: https://github.com/llvm/llvm-project/commit/c4a9a374749deb5f2a932a7d4ef9321be1b2ae5d
DIFF: https://github.com/llvm/llvm-project/commit/c4a9a374749deb5f2a932a7d4ef9321be1b2ae5d.diff
LOG: [OpenACC] device_type clause Sema for Compute constructs
device_type, also spelled as dtype, specifies the applicability of the
clauses following it, and takes a series of identifiers representing the
architectures it applies to. As we don't have a source for the valid
architectures yet, this patch just accepts all.
Semantically, this also limits the list of clauses that can be applied
after the device_type, so this implements that as well.
Added:
clang/test/SemaOpenACC/compute-construct-device_type-ast.cpp
clang/test/SemaOpenACC/compute-construct-device_type-clause.c
clang/test/SemaOpenACC/compute-construct-device_type-clause.cpp
Modified:
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenACCClauses.def
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/SemaOpenACC.h
clang/lib/AST/OpenACCClause.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/AST/TextNodeDumper.cpp
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/AST/ast-print-openacc-compute-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/tools/libclang/CIndex.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 3d0b1ab9d31e0..3998a3430ff95 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -17,6 +17,8 @@
#include "clang/AST/StmtIterator.h"
#include "clang/Basic/OpenACCKinds.h"
+#include <utility>
+
namespace clang {
/// This is the base type for all OpenACC Clauses.
class OpenACCClause {
@@ -75,6 +77,63 @@ class OpenACCClauseWithParams : public OpenACCClause {
}
};
+using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
+/// A 'device_type' or 'dtype' clause, takes a list of either an 'asterisk' or
+/// an identifier. The 'asterisk' means 'the rest'.
+class OpenACCDeviceTypeClause final
+ : public OpenACCClauseWithParams,
+ public llvm::TrailingObjects<OpenACCDeviceTypeClause,
+ DeviceTypeArgument> {
+ // Data stored in trailing objects as IdentifierInfo* /SourceLocation pairs. A
+ // nullptr IdentifierInfo* represents an asterisk.
+ unsigned NumArchs;
+ OpenACCDeviceTypeClause(OpenACCClauseKind K, SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<DeviceTypeArgument> Archs,
+ SourceLocation EndLoc)
+ : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc),
+ NumArchs(Archs.size()) {
+ assert(
+ (K == OpenACCClauseKind::DeviceType || K == OpenACCClauseKind::DType) &&
+ "Invalid clause kind for device-type");
+
+ assert(!llvm::any_of(Archs, [](const DeviceTypeArgument &Arg) {
+ return Arg.second.isInvalid();
+ }) && "Invalid SourceLocation for an argument");
+
+ assert(Archs.size() == 1 ||
+ !llvm::any_of(Archs,
+ [](const DeviceTypeArgument &Arg) {
+ return Arg.first == nullptr;
+ }) &&
+ "Only a single asterisk version is permitted, and must be the "
+ "only one");
+
+ std::uninitialized_copy(Archs.begin(), Archs.end(),
+ getTrailingObjects<DeviceTypeArgument>());
+ }
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::DType ||
+ C->getClauseKind() == OpenACCClauseKind::DeviceType;
+ }
+ bool hasAsterisk() const {
+ return getArchitectures().size() > 0 &&
+ getArchitectures()[0].first == nullptr;
+ }
+
+ ArrayRef<DeviceTypeArgument> getArchitectures() const {
+ return ArrayRef<DeviceTypeArgument>(
+ getTrailingObjects<DeviceTypeArgument>(), NumArchs);
+ }
+
+ static OpenACCDeviceTypeClause *
+ Create(const ASTContext &C, OpenACCClauseKind K, SourceLocation BeginLoc,
+ SourceLocation LParenLoc, ArrayRef<DeviceTypeArgument> Archs,
+ SourceLocation EndLoc);
+};
+
/// A 'default' clause, has the optional 'none' or 'present' argument.
class OpenACCDefaultClause : public OpenACCClauseWithParams {
friend class ASTReaderStmt;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 9e82130c93609..6100fba510059 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12344,4 +12344,8 @@ def warn_acc_deprecated_alias_name
def err_acc_var_not_pointer_type
: Error<"expected pointer in '%0' clause, type is %1">;
def note_acc_expected_pointer_var : Note<"expected variable of pointer type">;
+def err_acc_clause_after_device_type
+ : Error<"OpenACC clause '%0' may not follow a '%1' clause in a "
+ "compute construct">;
+
} // end of sema component.
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index afb7b30b7465c..7ecc51799468c 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -37,6 +37,8 @@ CLAUSE_ALIAS(PCreate, Create)
CLAUSE_ALIAS(PresentOrCreate, Create)
VISIT_CLAUSE(Default)
VISIT_CLAUSE(DevicePtr)
+VISIT_CLAUSE(DeviceType)
+CLAUSE_ALIAS(DType, DeviceType)
VISIT_CLAUSE(FirstPrivate)
VISIT_CLAUSE(If)
VISIT_CLAUSE(NoCreate)
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 61589fb7766f4..3910cba34a215 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3720,7 +3720,8 @@ class Parser : public CodeCompletionHandler {
SourceLocation Loc,
llvm::SmallVectorImpl<Expr *> &IntExprs);
/// Parses the 'device-type-list', which is a list of identifiers.
- bool ParseOpenACCDeviceTypeList();
+ bool ParseOpenACCDeviceTypeList(
+ llvm::SmallVector<std::pair<IdentifierInfo *, SourceLocation>> &Archs);
/// Parses the 'async-argument', which is an integral value with two
/// 'special' values that are likely negative (but come from Macros).
OpenACCIntExprParseResult ParseOpenACCAsyncArgument(OpenACCDirectiveKind DK,
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index e684ee6b2be12..f838fa97d33a2 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -26,6 +26,9 @@ class OpenACCClause;
class SemaOpenACC : public SemaBase {
public:
+ // Redeclaration of the version in OpenACCClause.h.
+ using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
+
/// A type to represent all the data for an OpenACC Clause that has been
/// parsed, but not yet created/semantically analyzed. This is effectively a
/// discriminated union on the 'Clause Kind', with all of the individual
@@ -60,8 +63,12 @@ class SemaOpenACC : public SemaBase {
SmallVector<Expr *> QueueIdExprs;
};
+ struct DeviceTypeDetails {
+ SmallVector<DeviceTypeArgument> Archs;
+ };
+
std::variant<std::monostate, DefaultDetails, ConditionDetails,
- IntExprDetails, VarListDetails, WaitDetails>
+ IntExprDetails, VarListDetails, WaitDetails, DeviceTypeDetails>
Details = std::monostate{};
public:
@@ -209,6 +216,13 @@ class SemaOpenACC : public SemaBase {
return std::get<VarListDetails>(Details).IsZero;
}
+ ArrayRef<DeviceTypeArgument> getDeviceTypeArchitectures() const {
+ assert((ClauseKind == OpenACCClauseKind::DeviceType ||
+ ClauseKind == OpenACCClauseKind::DType) &&
+ "Only 'device_type'/'dtype' has a device-type-arg list");
+ return std::get<DeviceTypeDetails>(Details).Archs;
+ }
+
void setLParenLoc(SourceLocation EndLoc) { LParenLoc = EndLoc; }
void setEndLoc(SourceLocation EndLoc) { ClauseRange.setEnd(EndLoc); }
@@ -326,6 +340,13 @@ class SemaOpenACC : public SemaBase {
"Parsed clause kind does not have a wait-details");
Details = WaitDetails{DevNum, QueuesLoc, std::move(IntExprs)};
}
+
+ void setDeviceTypeDetails(llvm::SmallVector<DeviceTypeArgument> &&Archs) {
+ assert((ClauseKind == OpenACCClauseKind::DeviceType ||
+ ClauseKind == OpenACCClauseKind::DType) &&
+ "Only 'device_type'/'dtype' has a device-type-arg list");
+ Details = DeviceTypeDetails{std::move(Archs)};
+ }
};
SemaOpenACC(Sema &S);
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index ee13437b97b48..f80ecc90d3963 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -18,7 +18,8 @@
using namespace clang;
bool OpenACCClauseWithParams::classof(const OpenACCClause *C) {
- return OpenACCClauseWithCondition::classof(C) ||
+ return OpenACCDeviceTypeClause::classof(C) ||
+ OpenACCClauseWithCondition::classof(C) ||
OpenACCClauseWithExprs::classof(C);
}
bool OpenACCClauseWithExprs::classof(const OpenACCClause *C) {
@@ -298,6 +299,17 @@ OpenACCCreateClause::Create(const ASTContext &C, OpenACCClauseKind Spelling,
VarList, EndLoc);
}
+OpenACCDeviceTypeClause *OpenACCDeviceTypeClause::Create(
+ const ASTContext &C, OpenACCClauseKind K, SourceLocation BeginLoc,
+ SourceLocation LParenLoc, ArrayRef<DeviceTypeArgument> Archs,
+ SourceLocation EndLoc) {
+ void *Mem =
+ C.Allocate(OpenACCDeviceTypeClause::totalSizeToAlloc<DeviceTypeArgument>(
+ Archs.size()));
+ return new (Mem)
+ OpenACCDeviceTypeClause(K, BeginLoc, LParenLoc, Archs, EndLoc);
+}
+
//===----------------------------------------------------------------------===//
// OpenACC clauses printing methods
//===----------------------------------------------------------------------===//
@@ -451,3 +463,17 @@ void OpenACCClausePrinter::VisitWaitClause(const OpenACCWaitClause &C) {
OS << ")";
}
}
+
+void OpenACCClausePrinter::VisitDeviceTypeClause(
+ const OpenACCDeviceTypeClause &C) {
+ OS << C.getClauseKind();
+ OS << "(";
+ llvm::interleaveComma(C.getArchitectures(), OS,
+ [&](const DeviceTypeArgument &Arch) {
+ if (Arch.first == nullptr)
+ OS << "*";
+ else
+ OS << Arch.first;
+ });
+ OS << ")";
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 8fb8940142eb0..caab4ab0ef160 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2585,6 +2585,9 @@ void OpenACCClauseProfiler::VisitWaitClause(const OpenACCWaitClause &Clause) {
for (auto *E : Clause.getQueueIdExprs())
Profiler.VisitStmt(E);
}
+/// Nothing to do here, there are no sub-statements.
+void OpenACCClauseProfiler::VisitDeviceTypeClause(
+ const OpenACCDeviceTypeClause &Clause) {}
} // namespace
void StmtProfiler::VisitOpenACCComputeConstruct(
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 12aa5858b7983..efcd74717a4e2 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -444,6 +444,19 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
if (cast<OpenACCWaitClause>(C)->hasQueuesTag())
OS << " has queues tag";
break;
+ case OpenACCClauseKind::DeviceType:
+ case OpenACCClauseKind::DType:
+ OS << "(";
+ llvm::interleaveComma(
+ cast<OpenACCDeviceTypeClause>(C)->getArchitectures(), OS,
+ [&](const DeviceTypeArgument &Arch) {
+ if (Arch.first == nullptr)
+ OS << "*";
+ else
+ OS << Arch.first->getName();
+ });
+ OS << ")";
+ break;
default:
// Nothing to do here.
break;
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 0e10632c83175..261c9cdc088b0 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -711,14 +711,15 @@ bool Parser::ParseOpenACCIntExprList(OpenACCDirectiveKind DK,
/// device_type( device-type-list )
///
/// The device_type clause may be abbreviated to dtype.
-bool Parser::ParseOpenACCDeviceTypeList() {
+bool Parser::ParseOpenACCDeviceTypeList(
+ llvm::SmallVector<std::pair<IdentifierInfo *, SourceLocation>> &Archs) {
if (expectIdentifierOrKeyword(*this)) {
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
Parser::StopBeforeMatch);
- return false;
+ return true;
}
- ConsumeToken();
+ Archs.emplace_back(getCurToken().getIdentifierInfo(), ConsumeToken());
while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
ExpectAndConsume(tok::comma);
@@ -726,9 +727,9 @@ bool Parser::ParseOpenACCDeviceTypeList() {
if (expectIdentifierOrKeyword(*this)) {
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
Parser::StopBeforeMatch);
- return false;
+ return true;
}
- ConsumeToken();
+ Archs.emplace_back(getCurToken().getIdentifierInfo(), ConsumeToken());
}
return false;
}
@@ -1021,16 +1022,20 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
break;
}
case OpenACCClauseKind::DType:
- case OpenACCClauseKind::DeviceType:
+ case OpenACCClauseKind::DeviceType: {
+ llvm::SmallVector<std::pair<IdentifierInfo *, SourceLocation>> Archs;
if (getCurToken().is(tok::star)) {
// FIXME: We want to mark that this is an 'everything else' type of
// device_type in Sema.
- ConsumeToken();
- } else if (ParseOpenACCDeviceTypeList()) {
+ ParsedClause.setDeviceTypeDetails({{nullptr, ConsumeToken()}});
+ } else if (!ParseOpenACCDeviceTypeList(Archs)) {
+ ParsedClause.setDeviceTypeDetails(std::move(Archs));
+ } else {
Parens.skipToEnd();
return OpenACCCanContinue();
}
break;
+ }
case OpenACCClauseKind::Tile:
if (ParseOpenACCSizeExprList()) {
Parens.skipToEnd();
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 656d30947a8d1..f174b2fa63c6a 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -255,6 +255,33 @@ bool checkAlreadyHasClauseOfKind(
return false;
}
+/// Implement check from OpenACC3.3: section 2.5.4:
+/// Only the async, wait, num_gangs, num_workers, and vector_length clauses may
+/// follow a device_type clause.
+bool checkValidAfterDeviceType(
+ SemaOpenACC &S, const OpenACCDeviceTypeClause &DeviceTypeClause,
+ const SemaOpenACC::OpenACCParsedClause &NewClause) {
+ // This is only a requirement on compute constructs so far, so this is fine
+ // otherwise.
+ if (!isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind()))
+ return false;
+ switch (NewClause.getClauseKind()) {
+ case OpenACCClauseKind::Async:
+ case OpenACCClauseKind::Wait:
+ case OpenACCClauseKind::NumGangs:
+ case OpenACCClauseKind::NumWorkers:
+ case OpenACCClauseKind::VectorLength:
+ case OpenACCClauseKind::DType:
+ case OpenACCClauseKind::DeviceType:
+ return false;
+ default:
+ S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type)
+ << NewClause.getClauseKind() << DeviceTypeClause.getClauseKind();
+ S.Diag(DeviceTypeClause.getBeginLoc(), diag::note_acc_previous_clause_here);
+ return true;
+ }
+}
+
} // namespace
SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {}
@@ -273,6 +300,17 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
return nullptr;
}
+ if (const auto *DevTypeClause =
+ llvm::find_if(ExistingClauses,
+ [&](const OpenACCClause *C) {
+ return isa<OpenACCDeviceTypeClause>(C);
+ });
+ DevTypeClause != ExistingClauses.end()) {
+ if (checkValidAfterDeviceType(
+ *this, *cast<OpenACCDeviceTypeClause>(*DevTypeClause), Clause))
+ return nullptr;
+ }
+
switch (Clause.getClauseKind()) {
case OpenACCClauseKind::Default: {
// Restrictions only properly implemented on 'compute' constructs, and
@@ -651,6 +689,23 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
Clause.getDevNumExpr(), Clause.getQueuesLoc(), Clause.getQueueIdExprs(),
Clause.getEndLoc());
}
+ case OpenACCClauseKind::DType:
+ case OpenACCClauseKind::DeviceType: {
+ // Restrictions only properly implemented on 'compute' constructs, and
+ // 'compute' constructs are the only construct that can do anything with
+ // this yet, so skip/treat as unimplemented in this case.
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+ break;
+
+ // TODO OpenACC: Once we get enough of the CodeGen implemented that we have
+ // a source for the list of valid architectures, we need to warn on unknown
+ // identifiers here.
+
+ return OpenACCDeviceTypeClause::Create(
+ getASTContext(), Clause.getClauseKind(), Clause.getBeginLoc(),
+ Clause.getLParenLoc(), Clause.getDeviceTypeArchitectures(),
+ Clause.getEndLoc());
+ }
default:
break;
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 126965088831d..ab26d1b1199ae 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11480,6 +11480,16 @@ void OpenACCClauseTransform<Derived>::VisitWaitClause(
ParsedClause.getQueuesLoc(), ParsedClause.getQueueIdExprs(),
ParsedClause.getEndLoc());
}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDeviceTypeClause(
+ const OpenACCDeviceTypeClause &C) {
+ // Nothing to transform here, just create a new version of 'C'.
+ NewClause = OpenACCDeviceTypeClause::Create(
+ Self.getSema().getASTContext(), C.getClauseKind(),
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ C.getArchitectures(), ParsedClause.getEndLoc());
+}
} // namespace
template <typename Derived>
OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 7627996d2c322..8f437a7c5f50a 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11905,6 +11905,21 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
DevNumExpr, QueuesLoc, QueueIdExprs,
EndLoc);
}
+ case OpenACCClauseKind::DeviceType:
+ case OpenACCClauseKind::DType: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<DeviceTypeArgument> Archs;
+ unsigned NumArchs = readInt();
+
+ for (unsigned I = 0; I < NumArchs; ++I) {
+ IdentifierInfo *Ident = readBool() ? readIdentifier() : nullptr;
+ SourceLocation Loc = readSourceLocation();
+ Archs.emplace_back(Ident, Loc);
+ }
+
+ return OpenACCDeviceTypeClause::Create(getContext(), ClauseKind, BeginLoc,
+ LParenLoc, Archs, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -11926,8 +11941,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
- case OpenACCClauseKind::DeviceType:
- case OpenACCClauseKind::DType:
case OpenACCClauseKind::Tile:
case OpenACCClauseKind::Gang:
case OpenACCClauseKind::Invalid:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 6154ead589d3e..7a9d392889bbd 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7933,6 +7933,19 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
writeOpenACCIntExprList(WC->getQueueIdExprs());
return;
}
+ case OpenACCClauseKind::DeviceType:
+ case OpenACCClauseKind::DType: {
+ const auto *DTC = cast<OpenACCDeviceTypeClause>(C);
+ writeSourceLocation(DTC->getLParenLoc());
+ writeUInt32(DTC->getArchitectures().size());
+ for (const DeviceTypeArgument &Arg : DTC->getArchitectures()) {
+ writeBool(Arg.first);
+ if (Arg.first)
+ AddIdentifierRef(Arg.first);
+ writeSourceLocation(Arg.second);
+ }
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -7954,8 +7967,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
- case OpenACCClauseKind::DeviceType:
- case OpenACCClauseKind::DType:
case OpenACCClauseKind::Tile:
case OpenACCClauseKind::Gang:
case OpenACCClauseKind::Invalid:
diff --git a/clang/test/AST/ast-print-openacc-compute-construct.cpp b/clang/test/AST/ast-print-openacc-compute-construct.cpp
index 0bfb90bcb5871..cdd9ab3377d01 100644
--- a/clang/test/AST/ast-print-openacc-compute-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-compute-construct.cpp
@@ -107,5 +107,28 @@ void foo() {
// CHECK: #pragma acc parallel wait(devnum: i : queues: *iPtr, i)
#pragma acc parallel wait(devnum:i:queues:*iPtr, i)
while(true);
+
+ bool SomeB;
+ struct SomeStruct{} SomeStructImpl;
+
+//#pragma acc parallel dtype(SomeB)
+#pragma acc parallel dtype(SomeB)
+ while(true);
+
+//#pragma acc parallel device_type(SomeStruct)
+#pragma acc parallel device_type(SomeStruct)
+ while(true);
+
+//#pragma acc parallel device_type(int)
+#pragma acc parallel device_type(int)
+ while(true);
+
+//#pragma acc parallel dtype(bool)
+#pragma acc parallel dtype(bool)
+ while(true);
+
+//#pragma acc parallel device_type (SomeStructImpl)
+#pragma acc parallel device_type (SomeStructImpl)
+ while(true);
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 51858b441e935..694f28b86ec9f 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1126,12 +1126,10 @@ void device_type() {
#pragma acc parallel dtype(
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'device_type' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel device_type()
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'dtype' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel dtype()
{}
@@ -1173,12 +1171,10 @@ void device_type() {
#pragma acc parallel dtype(ident, ident2
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'device_type' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel device_type(ident, ident2,)
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'dtype' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel dtype(ident, ident2,)
{}
@@ -1200,33 +1196,25 @@ void device_type() {
#pragma acc parallel dtype(*,ident)
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'device_type' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel device_type(ident, *)
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'dtype' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel dtype(ident, *)
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'device_type' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel device_type("foo", 54)
{}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC clause 'dtype' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc parallel dtype(31, "bar")
{}
- // expected-warning at +1{{OpenACC clause 'device_type' not yet implemented, clause ignored}}
#pragma acc parallel device_type(ident, auto, int, float)
{}
- // expected-warning at +1{{OpenACC clause 'dtype' not yet implemented, clause ignored}}
#pragma acc parallel dtype(ident, auto, int, float)
{}
- // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC clause 'dtype' not yet implemented, clause ignored}}
#pragma acc parallel device_type(ident, auto, int, float) dtype(ident, auto, int, float)
{}
}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-ast.cpp b/clang/test/SemaOpenACC/compute-construct-device_type-ast.cpp
new file mode 100644
index 0000000000000..8a2423f4f5427
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-ast.cpp
@@ -0,0 +1,105 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+struct SomeS{};
+void NormalUses() {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK-NEXT: CompoundStmt
+
+ SomeS SomeImpl;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} SomeImpl 'SomeS'
+ // CHECK-NEXT: CXXConstructExpr
+ bool SomeVar;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} SomeVar 'bool'
+
+#pragma acc parallel device_type(SomeS) dtype(SomeImpl)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(SomeS)
+ // CHECK-NEXT: dtype(SomeImpl)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+#pragma acc parallel device_type(SomeVar) dtype(int)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(SomeVar)
+ // CHECK-NEXT: dtype(int)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+#pragma acc parallel device_type(private) dtype(struct)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(private)
+ // CHECK-NEXT: dtype(struct)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+#pragma acc parallel device_type(private) dtype(class)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(private)
+ // CHECK-NEXT: dtype(class)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+#pragma acc parallel device_type(float) dtype(*)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(float)
+ // CHECK-NEXT: dtype(*)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+#pragma acc parallel device_type(float, int) dtype(*)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(float, int)
+ // CHECK-NEXT: dtype(*)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+}
+
+template<typename T>
+void TemplUses() {
+ // CHECK-NEXT: FunctionTemplateDecl{{.*}}TemplUses
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
+ // CHECK-NEXT: FunctionDecl{{.*}}TemplUses
+ // CHECK-NEXT: CompoundStmt
+#pragma acc parallel device_type(T) dtype(T)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(T)
+ // CHECK-NEXT: dtype(T)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+
+
+ // Instantiations
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: device_type(T)
+ // CHECK-NEXT: dtype(T)
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+}
+
+void Inst() {
+ TemplUses<int>();
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
new file mode 100644
index 0000000000000..15c9cf396c80c
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -0,0 +1,221 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+#define MACRO +FOO
+
+void uses() {
+ typedef struct S{} STy;
+ STy SImpl;
+
+#pragma acc parallel device_type(I)
+ while(1);
+#pragma acc serial device_type(S) dtype(STy)
+ while(1);
+#pragma acc kernels dtype(SImpl)
+ while(1);
+#pragma acc kernels dtype(int) device_type(*)
+ while(1);
+#pragma acc kernels dtype(true) device_type(false)
+ while(1);
+
+ // expected-error at +1{{expected identifier}}
+#pragma acc kernels dtype(int, *)
+ while(1);
+
+#pragma acc parallel device_type(I, int)
+ while(1);
+ // expected-error at +2{{expected ','}}
+ // expected-error at +1{{expected identifier}}
+#pragma acc kernels dtype(int{})
+ while(1);
+ // expected-error at +1{{expected identifier}}
+#pragma acc kernels dtype(5)
+ while(1);
+ // expected-error at +1{{expected identifier}}
+#pragma acc kernels dtype(MACRO)
+ while(1);
+
+
+ // Only 'async', 'wait', num_gangs', 'num_workers', 'vector_length' allowed after 'device_type'.
+
+ // expected-error at +2{{OpenACC clause 'finalize' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) finalize
+ while(1);
+ // expected-error at +2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) if_present
+ while(1);
+ // expected-error at +2{{OpenACC clause 'seq' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) seq
+ while(1);
+ // expected-error at +2{{OpenACC clause 'independent' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) independent
+ while(1);
+ // expected-error at +2{{OpenACC clause 'auto' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) auto
+ while(1);
+ // expected-error at +2{{OpenACC clause 'worker' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) worker
+ while(1);
+ // expected-error at +2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) nohost
+ while(1);
+ // expected-error at +2{{OpenACC clause 'default' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) default(none)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'if' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) if(1)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'self' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) self
+ while(1);
+
+ int Var;
+ int *VarPtr;
+ // expected-error at +2{{OpenACC clause 'copy' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) copy(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'pcopy' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) pcopy(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'present_or_copy' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) present_or_copy(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'use_device' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) use_device(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'attach' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) attach(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'delete' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) delete(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) detach(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'device' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) device(VarPtr)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'deviceptr' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) deviceptr(VarPtr)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) device_resident(VarPtr)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel device_type(*) firstprivate(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'host' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) host(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) link(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) no_create(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'present' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) present(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'private' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel device_type(*) private(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'copyout' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) copyout(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'pcopyout' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) pcopyout(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'present_or_copyout' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) present_or_copyout(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'copyin' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) copyin(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'pcopyin' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) pcopyin(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'present_or_copyin' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) present_or_copyin(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'create' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) create(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'pcreate' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) pcreate(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'present_or_create' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) present_or_create(Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'reduction' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) reduction(+:Var)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'collapse' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) collapse(1)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) bind(Var)
+ while(1);
+#pragma acc kernels device_type(*) vector_length(1)
+ while(1);
+#pragma acc kernels device_type(*) num_gangs(1)
+ while(1);
+#pragma acc kernels device_type(*) num_workers(1)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'device_num' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) device_num(1)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'default_async' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) default_async(1)
+ while(1);
+#pragma acc kernels device_type(*) async
+ while(1);
+ // expected-error at +2{{OpenACC clause 'tile' may not follow a 'device_type' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels device_type(*) tile(Var, 1)
+ while(1);
+ // expected-error at +2{{OpenACC clause 'gang' may not follow a 'dtype' clause in a compute construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels dtype(*) gang
+ while(1);
+#pragma acc kernels device_type(*) wait
+ while(1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.cpp b/clang/test/SemaOpenACC/compute-construct-device_type-clause.cpp
new file mode 100644
index 0000000000000..ed40e8bbceae7
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.cpp
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+template<typename T>
+void TemplUses() {
+#pragma acc parallel device_type(I)
+ while(true);
+#pragma acc parallel dtype(*)
+ while(true);
+#pragma acc parallel device_type(class)
+ while(true);
+#pragma acc parallel device_type(private)
+ while(true);
+#pragma acc parallel device_type(bool)
+ while(true);
+#pragma acc kernels dtype(true) device_type(false)
+ while(true);
+ // expected-error at +2{{expected ','}}
+ // expected-error at +1{{expected identifier}}
+#pragma acc parallel device_type(T::value)
+ while(true);
+}
+
+void Inst() {
+ TemplUses<int>(); // #INST
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index ae6659fe95e89..8b9417f985b50 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2857,6 +2857,8 @@ void OpenACCClauseEnqueue::VisitWaitClause(const OpenACCWaitClause &C) {
for (Expr *QE : C.getQueueIdExprs())
Visitor.AddStmt(QE);
}
+void OpenACCClauseEnqueue::VisitDeviceTypeClause(
+ const OpenACCDeviceTypeClause &C) {}
} // namespace
void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {
More information about the cfe-commits
mailing list