Author: erichkeane Date: 2024-05-13T10:29:43-07:00 New Revision: 8ef2011b2cd3a8fc2ef8d6ea0facb1a39a0dd621
URL: https://github.com/llvm/llvm-project/commit/8ef2011b2cd3a8fc2ef8d6ea0facb1a39a0dd621 DIFF: https://github.com/llvm/llvm-project/commit/8ef2011b2cd3a8fc2ef8d6ea0facb1a39a0dd621.diff LOG: Reapply "[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. This reverts commit 06f04b2e27f2586d3db2204ed4e54f8b78fea74e. This reapplies commit c4a9a374749deb5f2a932a7d4ef9321be1b2ae5d. The build failures were caused by the patch depending on the order of evaluation of arguments to a function. This reapplication separates out the capture of one of the values. 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..607a2b9d65367 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..5db3036b00030 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -711,14 +711,16 @@ 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(); + IdentifierInfo *Ident = getCurToken().getIdentifierInfo(); + Archs.emplace_back(Ident, ConsumeToken()); while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) { ExpectAndConsume(tok::comma); @@ -726,9 +728,10 @@ bool Parser::ParseOpenACCDeviceTypeList() { if (expectIdentifierOrKeyword(*this)) { SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, Parser::StopBeforeMatch); - return false; + return true; } - ConsumeToken(); + Ident = getCurToken().getIdentifierInfo(); + Archs.emplace_back(Ident, ConsumeToken()); } return false; } @@ -1021,16 +1024,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@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'device_type' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel device_type() {} - // expected-error@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'dtype' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel dtype() {} @@ -1173,12 +1171,10 @@ void device_type() { #pragma acc parallel dtype(ident, ident2 {} - // expected-error@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'device_type' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel device_type(ident, ident2,) {} - // expected-error@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'dtype' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel dtype(ident, ident2,) {} @@ -1200,33 +1196,25 @@ void device_type() { #pragma acc parallel dtype(*,ident) {} - // expected-error@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'device_type' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel device_type(ident, *) {} - // expected-error@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'dtype' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel dtype(ident, *) {} - // expected-error@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'device_type' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel device_type("foo", 54) {} - // expected-error@+2{{expected identifier}} - // expected-warning@+1{{OpenACC clause 'dtype' not yet implemented, clause ignored}} + // expected-error@+1{{expected identifier}} #pragma acc parallel dtype(31, "bar") {} - // expected-warning@+1{{OpenACC clause 'device_type' not yet implemented, clause ignored}} #pragma acc parallel device_type(ident, auto, int, float) {} - // expected-warning@+1{{OpenACC clause 'dtype' not yet implemented, clause ignored}} #pragma acc parallel dtype(ident, auto, int, float) {} - // expected-warning@+2{{OpenACC clause 'device_type' not yet implemented, clause ignored}} - // expected-warning@+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@+1{{expected identifier}} +#pragma acc kernels dtype(int, *) + while(1); + +#pragma acc parallel device_type(I, int) + while(1); + // expected-error@+2{{expected ','}} + // expected-error@+1{{expected identifier}} +#pragma acc kernels dtype(int{}) + while(1); + // expected-error@+1{{expected identifier}} +#pragma acc kernels dtype(5) + while(1); + // expected-error@+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@+2{{OpenACC clause 'finalize' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) finalize + while(1); + // expected-error@+2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) if_present + while(1); + // expected-error@+2{{OpenACC clause 'seq' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) seq + while(1); + // expected-error@+2{{OpenACC clause 'independent' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) independent + while(1); + // expected-error@+2{{OpenACC clause 'auto' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) auto + while(1); + // expected-error@+2{{OpenACC clause 'worker' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) worker + while(1); + // expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) nohost + while(1); + // expected-error@+2{{OpenACC clause 'default' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) default(none) + while(1); + // expected-error@+2{{OpenACC clause 'if' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) if(1) + while(1); + // expected-error@+2{{OpenACC clause 'self' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) self + while(1); + + int Var; + int *VarPtr; + // expected-error@+2{{OpenACC clause 'copy' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) copy(Var) + while(1); + // expected-error@+2{{OpenACC clause 'pcopy' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) pcopy(Var) + while(1); + // expected-error@+2{{OpenACC clause 'present_or_copy' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) present_or_copy(Var) + while(1); + // expected-error@+2{{OpenACC clause 'use_device' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) use_device(Var) + while(1); + // expected-error@+2{{OpenACC clause 'attach' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) attach(Var) + while(1); + // expected-error@+2{{OpenACC clause 'delete' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) delete(Var) + while(1); + // expected-error@+2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) detach(Var) + while(1); + // expected-error@+2{{OpenACC clause 'device' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) device(VarPtr) + while(1); + // expected-error@+2{{OpenACC clause 'deviceptr' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) deviceptr(VarPtr) + while(1); + // expected-error@+2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) device_resident(VarPtr) + while(1); + // expected-error@+2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel device_type(*) firstprivate(Var) + while(1); + // expected-error@+2{{OpenACC clause 'host' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) host(Var) + while(1); + // expected-error@+2{{OpenACC clause 'link' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) link(Var) + while(1); + // expected-error@+2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) no_create(Var) + while(1); + // expected-error@+2{{OpenACC clause 'present' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) present(Var) + while(1); + // expected-error@+2{{OpenACC clause 'private' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel device_type(*) private(Var) + while(1); + // expected-error@+2{{OpenACC clause 'copyout' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) copyout(Var) + while(1); + // expected-error@+2{{OpenACC clause 'pcopyout' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) pcopyout(Var) + while(1); + // expected-error@+2{{OpenACC clause 'present_or_copyout' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) present_or_copyout(Var) + while(1); + // expected-error@+2{{OpenACC clause 'copyin' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) copyin(Var) + while(1); + // expected-error@+2{{OpenACC clause 'pcopyin' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) pcopyin(Var) + while(1); + // expected-error@+2{{OpenACC clause 'present_or_copyin' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) present_or_copyin(Var) + while(1); + // expected-error@+2{{OpenACC clause 'create' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) create(Var) + while(1); + // expected-error@+2{{OpenACC clause 'pcreate' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) pcreate(Var) + while(1); + // expected-error@+2{{OpenACC clause 'present_or_create' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) present_or_create(Var) + while(1); + // expected-error@+2{{OpenACC clause 'reduction' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) reduction(+:Var) + while(1); + // expected-error@+2{{OpenACC clause 'collapse' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) collapse(1) + while(1); + // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+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@+2{{OpenACC clause 'device_num' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) device_num(1) + while(1); + // expected-error@+2{{OpenACC clause 'default_async' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+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@+2{{OpenACC clause 'tile' may not follow a 'device_type' clause in a compute construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels device_type(*) tile(Var, 1) + while(1); + // expected-error@+2{{OpenACC clause 'gang' may not follow a 'dtype' clause in a compute construct}} + // expected-note@+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@+2{{expected ','}} + // expected-error@+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) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
