Author: erichkeane Date: 2024-05-06T09:29:04-07:00 New Revision: 48c8a5791ae71c96661479f684459b7b9427a22d
URL: https://github.com/llvm/llvm-project/commit/48c8a5791ae71c96661479f684459b7b9427a22d DIFF: https://github.com/llvm/llvm-project/commit/48c8a5791ae71c96661479f684459b7b9427a22d.diff LOG: [OpenACC] Implement 'deviceptr' and 'attach' sema for compute constructs These two are very similar to the other 'var-list' variants, except they require that the type of the variable be a pointer. This patch implements that restriction. Added: clang/test/SemaOpenACC/compute-construct-attach-clause.c clang/test/SemaOpenACC/compute-construct-attach-clause.cpp clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp Modified: clang/include/clang/AST/OpenACCClause.h clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Basic/OpenACCClauses.def 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/test/SemaOpenACC/compute-construct-varlist-ast.cpp clang/tools/libclang/CIndex.cpp Removed: ################################################################################ diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index abbd7c3989bcca..ec6b4aebcb9f4c 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -313,6 +313,44 @@ class OpenACCFirstPrivateClause final ArrayRef<Expr *> VarList, SourceLocation EndLoc); }; +class OpenACCDevicePtrClause final + : public OpenACCClauseWithVarList, + public llvm::TrailingObjects<OpenACCDevicePtrClause, Expr *> { + + OpenACCDevicePtrClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<Expr *> VarList, SourceLocation EndLoc) + : OpenACCClauseWithVarList(OpenACCClauseKind::DevicePtr, BeginLoc, + LParenLoc, EndLoc) { + std::uninitialized_copy(VarList.begin(), VarList.end(), + getTrailingObjects<Expr *>()); + setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size())); + } + +public: + static OpenACCDevicePtrClause * + Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<Expr *> VarList, SourceLocation EndLoc); +}; + +class OpenACCAttachClause final + : public OpenACCClauseWithVarList, + public llvm::TrailingObjects<OpenACCAttachClause, Expr *> { + + OpenACCAttachClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<Expr *> VarList, SourceLocation EndLoc) + : OpenACCClauseWithVarList(OpenACCClauseKind::Attach, BeginLoc, LParenLoc, + EndLoc) { + std::uninitialized_copy(VarList.begin(), VarList.end(), + getTrailingObjects<Expr *>()); + setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size())); + } + +public: + static OpenACCAttachClause * + Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<Expr *> VarList, SourceLocation EndLoc); +}; + class OpenACCNoCreateClause final : public OpenACCClauseWithVarList, public llvm::TrailingObjects<OpenACCNoCreateClause, Expr *> { diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 6c4d92790afc51..9a0bae9c216de9 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12337,4 +12337,7 @@ def warn_acc_deprecated_alias_name : Warning<"OpenACC clause name '%0' is a deprecated clause name and is " "now an alias for '%1'">, InGroup<DiagGroup<"openacc-deprecated-clause-alias">>; +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">; } // end of sema component. diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def index 06c7a379b22261..c92e5eb1e1b634 100644 --- a/clang/include/clang/Basic/OpenACCClauses.def +++ b/clang/include/clang/Basic/OpenACCClauses.def @@ -21,6 +21,7 @@ #define CLAUSE_ALIAS(ALIAS_NAME, CLAUSE_NAME) #endif +VISIT_CLAUSE(Attach) VISIT_CLAUSE(Copy) CLAUSE_ALIAS(PCopy, Copy) CLAUSE_ALIAS(PresentOrCopy, Copy) @@ -34,6 +35,7 @@ VISIT_CLAUSE(Create) CLAUSE_ALIAS(PCreate, Create) CLAUSE_ALIAS(PresentOrCreate, Create) VISIT_CLAUSE(Default) +VISIT_CLAUSE(DevicePtr) VISIT_CLAUSE(FirstPrivate) VISIT_CLAUSE(If) VISIT_CLAUSE(NoCreate) diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 5f77ec90d0b650..32d94ee8f33fed 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -134,6 +134,8 @@ class SemaOpenACC : public SemaBase { ClauseKind == OpenACCClauseKind::Create || ClauseKind == OpenACCClauseKind::PCreate || ClauseKind == OpenACCClauseKind::PresentOrCreate || + ClauseKind == OpenACCClauseKind::Attach || + ClauseKind == OpenACCClauseKind::DevicePtr || ClauseKind == OpenACCClauseKind::FirstPrivate) && "Parsed clause kind does not have a var-list"); return std::get<VarListDetails>(Details).VarList; @@ -217,6 +219,8 @@ class SemaOpenACC : public SemaBase { ClauseKind == OpenACCClauseKind::Create || ClauseKind == OpenACCClauseKind::PCreate || ClauseKind == OpenACCClauseKind::PresentOrCreate || + ClauseKind == OpenACCClauseKind::Attach || + ClauseKind == OpenACCClauseKind::DevicePtr || ClauseKind == OpenACCClauseKind::FirstPrivate) && "Parsed clause kind does not have a var-list"); assert((!IsReadOnly || ClauseKind == OpenACCClauseKind::CopyIn || @@ -251,6 +255,8 @@ class SemaOpenACC : public SemaBase { ClauseKind == OpenACCClauseKind::Create || ClauseKind == OpenACCClauseKind::PCreate || ClauseKind == OpenACCClauseKind::PresentOrCreate || + ClauseKind == OpenACCClauseKind::Attach || + ClauseKind == OpenACCClauseKind::DevicePtr || ClauseKind == OpenACCClauseKind::FirstPrivate) && "Parsed clause kind does not have a var-list"); assert((!IsReadOnly || ClauseKind == OpenACCClauseKind::CopyIn || @@ -315,6 +321,10 @@ class SemaOpenACC : public SemaBase { /// declaration reference to a variable of the correct type. ExprResult ActOnVar(Expr *VarExpr); + /// Called to check the 'var' type is a variable of pointer type, necessary + /// for 'deviceptr' and 'attach' clauses. Returns true on success. + bool CheckVarIsPointerType(OpenACCClauseKind ClauseKind, Expr *VarExpr); + /// Checks and creates an Array Section used in an OpenACC construct/clause. ExprResult ActOnArraySectionExpr(Expr *Base, SourceLocation LBLoc, Expr *LowerBound, diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index f682100e35d37b..c1affa97b781ca 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -156,6 +156,26 @@ OpenACCFirstPrivateClause *OpenACCFirstPrivateClause::Create( OpenACCFirstPrivateClause(BeginLoc, LParenLoc, VarList, EndLoc); } +OpenACCAttachClause *OpenACCAttachClause::Create(const ASTContext &C, + SourceLocation BeginLoc, + SourceLocation LParenLoc, + ArrayRef<Expr *> VarList, + SourceLocation EndLoc) { + void *Mem = + C.Allocate(OpenACCAttachClause::totalSizeToAlloc<Expr *>(VarList.size())); + return new (Mem) OpenACCAttachClause(BeginLoc, LParenLoc, VarList, EndLoc); +} + +OpenACCDevicePtrClause *OpenACCDevicePtrClause::Create(const ASTContext &C, + SourceLocation BeginLoc, + SourceLocation LParenLoc, + ArrayRef<Expr *> VarList, + SourceLocation EndLoc) { + void *Mem = C.Allocate( + OpenACCDevicePtrClause::totalSizeToAlloc<Expr *>(VarList.size())); + return new (Mem) OpenACCDevicePtrClause(BeginLoc, LParenLoc, VarList, EndLoc); +} + OpenACCNoCreateClause *OpenACCNoCreateClause::Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, @@ -282,6 +302,21 @@ void OpenACCClausePrinter::VisitFirstPrivateClause( OS << ")"; } +void OpenACCClausePrinter::VisitAttachClause(const OpenACCAttachClause &C) { + OS << "attach("; + llvm::interleaveComma(C.getVarList(), OS, + [&](const Expr *E) { printExpr(E); }); + OS << ")"; +} + +void OpenACCClausePrinter::VisitDevicePtrClause( + const OpenACCDevicePtrClause &C) { + OS << "deviceptr("; + llvm::interleaveComma(C.getVarList(), OS, + [&](const Expr *E) { printExpr(E); }); + OS << ")"; +} + void OpenACCClausePrinter::VisitNoCreateClause(const OpenACCNoCreateClause &C) { OS << "no_create("; llvm::interleaveComma(C.getVarList(), OS, diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index b97c351f83dbf4..11d3f3d4cec444 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2543,6 +2543,18 @@ void OpenACCClauseProfiler::VisitFirstPrivateClause( Profiler.VisitStmt(E); } +void OpenACCClauseProfiler::VisitAttachClause( + const OpenACCAttachClause &Clause) { + for (auto *E : Clause.getVarList()) + Profiler.VisitStmt(E); +} + +void OpenACCClauseProfiler::VisitDevicePtrClause( + const OpenACCDevicePtrClause &Clause) { + for (auto *E : Clause.getVarList()) + Profiler.VisitStmt(E); +} + void OpenACCClauseProfiler::VisitNoCreateClause( const OpenACCNoCreateClause &Clause) { for (auto *E : Clause.getVarList()) diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 5aeeb80e8fe710..21167ca56e5947 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -397,10 +397,12 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { case OpenACCClauseKind::Default: OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')'; break; + case OpenACCClauseKind::Attach: case OpenACCClauseKind::Copy: case OpenACCClauseKind::PCopy: case OpenACCClauseKind::PresentOrCopy: case OpenACCClauseKind::If: + case OpenACCClauseKind::DevicePtr: case OpenACCClauseKind::FirstPrivate: case OpenACCClauseKind::NoCreate: case OpenACCClauseKind::NumGangs: diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index c90ae8806f0868..b4b81e2ba13ea6 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -945,17 +945,21 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( // make sure we get the right diff erentiator. assert(DirKind == OpenACCDirectiveKind::Update); [[fallthrough]]; - case OpenACCClauseKind::Attach: case OpenACCClauseKind::Delete: case OpenACCClauseKind::Detach: case OpenACCClauseKind::Device: case OpenACCClauseKind::DeviceResident: - case OpenACCClauseKind::DevicePtr: case OpenACCClauseKind::Host: case OpenACCClauseKind::Link: case OpenACCClauseKind::UseDevice: ParseOpenACCVarList(); break; + case OpenACCClauseKind::Attach: + case OpenACCClauseKind::DevicePtr: + // TODO: ERICH: Figure out how to limit to just ptrs? + ParsedClause.setVarListDetails(ParseOpenACCVarList(), + /*IsReadOnly=*/false, /*IsZero=*/false); + break; case OpenACCClauseKind::Copy: case OpenACCClauseKind::PCopy: case OpenACCClauseKind::PresentOrCopy: diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 4e1a8c1277c5f0..8cf829cf215b54 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -170,6 +170,34 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, default: return false; } + case OpenACCClauseKind::Attach: + switch (DirectiveKind) { + case OpenACCDirectiveKind::Parallel: + case OpenACCDirectiveKind::Serial: + case OpenACCDirectiveKind::Kernels: + case OpenACCDirectiveKind::Data: + case OpenACCDirectiveKind::EnterData: + case OpenACCDirectiveKind::ParallelLoop: + case OpenACCDirectiveKind::SerialLoop: + case OpenACCDirectiveKind::KernelsLoop: + return true; + default: + return false; + } + case OpenACCClauseKind::DevicePtr: + switch (DirectiveKind) { + case OpenACCDirectiveKind::Parallel: + case OpenACCDirectiveKind::Serial: + case OpenACCDirectiveKind::Kernels: + case OpenACCDirectiveKind::Data: + case OpenACCDirectiveKind::Declare: + case OpenACCDirectiveKind::ParallelLoop: + case OpenACCDirectiveKind::SerialLoop: + case OpenACCDirectiveKind::KernelsLoop: + return true; + default: + return false; + } default: // Do nothing so we can go to the 'unimplemented' diagnostic instead. return true; @@ -513,6 +541,48 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses, Clause.getLParenLoc(), Clause.isZero(), Clause.getVarList(), Clause.getEndLoc()); } + case OpenACCClauseKind::Attach: { + // 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; + + // ActOnVar ensured that everything is a valid variable reference, but we + // still have to make sure it is a pointer type. + llvm::SmallVector<Expr *> VarList{Clause.getVarList().begin(), + Clause.getVarList().end()}; + VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) { + return CheckVarIsPointerType(OpenACCClauseKind::Attach, E); + }), VarList.end()); + Clause.setVarListDetails(VarList, + /*IsReadOnly=*/false, /*IsZero=*/false); + + return OpenACCAttachClause::Create(getASTContext(), Clause.getBeginLoc(), + Clause.getLParenLoc(), + Clause.getVarList(), Clause.getEndLoc()); + } + case OpenACCClauseKind::DevicePtr: { + // 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; + + // ActOnVar ensured that everything is a valid variable reference, but we + // still have to make sure it is a pointer type. + llvm::SmallVector<Expr *> VarList{Clause.getVarList().begin(), + Clause.getVarList().end()}; + VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) { + return CheckVarIsPointerType(OpenACCClauseKind::DevicePtr, E); + }), VarList.end()); + Clause.setVarListDetails(VarList, + /*IsReadOnly=*/false, /*IsZero=*/false); + + return OpenACCDevicePtrClause::Create( + getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(), + Clause.getVarList(), Clause.getEndLoc()); + } default: break; } @@ -641,6 +711,36 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK, return IntExpr; } +bool SemaOpenACC::CheckVarIsPointerType(OpenACCClauseKind ClauseKind, + Expr *VarExpr) { + // We already know that VarExpr is a proper reference to a variable, so we + // should be able to just take the type of the expression to get the type of + // the referenced variable. + + // We've already seen an error, don't diagnose anything else. + if (!VarExpr || VarExpr->containsErrors()) + return false; + + if (isa<ArraySectionExpr>(VarExpr->IgnoreParenImpCasts()) || + VarExpr->hasPlaceholderType(BuiltinType::ArraySection)) { + Diag(VarExpr->getExprLoc(), diag::err_array_section_use) << /*OpenACC=*/0; + Diag(VarExpr->getExprLoc(), diag::note_acc_expected_pointer_var); + return true; + } + + QualType Ty = VarExpr->getType(); + Ty = Ty.getNonReferenceType().getUnqualifiedType(); + + // Nothing we can do if this is a dependent type. + if (Ty->isDependentType()) + return false; + + if (!Ty->isPointerType()) + return Diag(VarExpr->getExprLoc(), diag::err_acc_var_not_pointer_type) + << ClauseKind << Ty; + return false; +} + ExprResult SemaOpenACC::ActOnVar(Expr *VarExpr) { // We still need to retain the array subscript/subarray exprs, so work on a // copy. diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index edf55703874688..a4ca8b5771a9f2 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11318,6 +11318,43 @@ void OpenACCClauseTransform<Derived>::VisitCreateClause( ParsedClause.isZero(), ParsedClause.getVarList(), ParsedClause.getEndLoc()); } +template <typename Derived> +void OpenACCClauseTransform<Derived>::VisitAttachClause( + const OpenACCAttachClause &C) { + llvm::SmallVector<Expr *> VarList = VisitVarList(C.getVarList()); + + // Ensure each var is a pointer type. + VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) { + return Self.getSema().OpenACC().CheckVarIsPointerType( + OpenACCClauseKind::Attach, E); + }), VarList.end()); + + ParsedClause.setVarListDetails(VarList, + /*IsReadOnly=*/false, /*IsZero=*/false); + NewClause = OpenACCAttachClause::Create( + Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), ParsedClause.getVarList(), + ParsedClause.getEndLoc()); +} + +template <typename Derived> +void OpenACCClauseTransform<Derived>::VisitDevicePtrClause( + const OpenACCDevicePtrClause &C) { + llvm::SmallVector<Expr *> VarList = VisitVarList(C.getVarList()); + + // Ensure each var is a pointer type. + VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) { + return Self.getSema().OpenACC().CheckVarIsPointerType( + OpenACCClauseKind::DevicePtr, E); + }), VarList.end()); + + ParsedClause.setVarListDetails(VarList, + /*IsReadOnly=*/false, /*IsZero=*/false); + NewClause = OpenACCDevicePtrClause::Create( + Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), ParsedClause.getVarList(), + ParsedClause.getEndLoc()); +} template <typename Derived> void OpenACCClauseTransform<Derived>::VisitNumWorkersClause( diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 4050d66d45170c..81b78edd9c6cfe 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -11822,6 +11822,18 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { return OpenACCFirstPrivateClause::Create(getContext(), BeginLoc, LParenLoc, VarList, EndLoc); } + case OpenACCClauseKind::Attach: { + SourceLocation LParenLoc = readSourceLocation(); + llvm::SmallVector<Expr *> VarList = readOpenACCVarList(); + return OpenACCAttachClause::Create(getContext(), BeginLoc, LParenLoc, + VarList, EndLoc); + } + case OpenACCClauseKind::DevicePtr: { + SourceLocation LParenLoc = readSourceLocation(); + llvm::SmallVector<Expr *> VarList = readOpenACCVarList(); + return OpenACCDevicePtrClause::Create(getContext(), BeginLoc, LParenLoc, + VarList, EndLoc); + } case OpenACCClauseKind::NoCreate: { SourceLocation LParenLoc = readSourceLocation(); llvm::SmallVector<Expr *> VarList = readOpenACCVarList(); @@ -11879,11 +11891,9 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { case OpenACCClauseKind::Vector: case OpenACCClauseKind::NoHost: case OpenACCClauseKind::UseDevice: - case OpenACCClauseKind::Attach: case OpenACCClauseKind::Delete: case OpenACCClauseKind::Detach: case OpenACCClauseKind::Device: - case OpenACCClauseKind::DevicePtr: case OpenACCClauseKind::DeviceResident: case OpenACCClauseKind::Host: case OpenACCClauseKind::Link: diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index cf77b4c0df2bbb..8a0116fa893247 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -7820,6 +7820,18 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { writeOpenACCVarList(FPC); return; } + case OpenACCClauseKind::Attach: { + const auto *AC = cast<OpenACCAttachClause>(C); + writeSourceLocation(AC->getLParenLoc()); + writeOpenACCVarList(AC); + return; + } + case OpenACCClauseKind::DevicePtr: { + const auto *DPC = cast<OpenACCDevicePtrClause>(C); + writeSourceLocation(DPC->getLParenLoc()); + writeOpenACCVarList(DPC); + return; + } case OpenACCClauseKind::NoCreate: { const auto *NCC = cast<OpenACCNoCreateClause>(C); writeSourceLocation(NCC->getLParenLoc()); @@ -7877,11 +7889,9 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { case OpenACCClauseKind::Vector: case OpenACCClauseKind::NoHost: case OpenACCClauseKind::UseDevice: - case OpenACCClauseKind::Attach: case OpenACCClauseKind::Delete: case OpenACCClauseKind::Detach: case OpenACCClauseKind::Device: - case OpenACCClauseKind::DevicePtr: case OpenACCClauseKind::DeviceResident: case OpenACCClauseKind::Host: case OpenACCClauseKind::Link: diff --git a/clang/test/AST/ast-print-openacc-compute-construct.cpp b/clang/test/AST/ast-print-openacc-compute-construct.cpp index 112f328f5cb9ce..1ee1e15bdfc3ac 100644 --- a/clang/test/AST/ast-print-openacc-compute-construct.cpp +++ b/clang/test/AST/ast-print-openacc-compute-construct.cpp @@ -2,7 +2,9 @@ void foo() { int i; + int *iPtr; float array[5]; + float *arrayPtr[5]; // CHECK: #pragma acc parallel default(none) // CHECK-NEXT: while (true) #pragma acc parallel default(none) @@ -65,5 +67,13 @@ void foo() { // CHECK: #pragma acc parallel create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2]) #pragma acc parallel create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2]) while(true); + + // CHECK: #pragma acc serial attach(iPtr, arrayPtr[0]) +#pragma acc serial attach(iPtr, arrayPtr[0]) + while(true); + + // CHECK: #pragma acc kernels deviceptr(iPtr, arrayPtr[0]) +#pragma acc kernels deviceptr(iPtr, arrayPtr[0]) + while(true); } diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index 65247e4db63efa..035b7ab4c1f40f 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -548,26 +548,30 @@ void VarListClauses() { #pragma acc serial present(s.array[s.value : 5], s.value), seq for(;;){} - // expected-error@+3{{expected ','}} - // expected-warning@+2{{OpenACC clause 'deviceptr' not yet implemented, clause ignored}} + + void *IsPointer; + // expected-error@+5{{expected ','}} + // expected-error@+4{{expected pointer in 'deviceptr' clause, type is 'char'}} + // expected-error@+3{{OpenACC sub-array is not allowed here}} + // expected-note@+2{{expected variable of pointer type}} // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}} #pragma acc serial deviceptr(s.array[s.value] s.array[s.value :5] ), seq for(;;){} - // expected-warning@+2{{OpenACC clause 'deviceptr' not yet implemented, clause ignored}} // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}} -#pragma acc serial deviceptr(s.array[s.value : 5], s.value), seq +#pragma acc serial deviceptr(IsPointer), seq for(;;){} - // expected-error@+3{{expected ','}} - // expected-warning@+2{{OpenACC clause 'attach' not yet implemented, clause ignored}} + // expected-error@+5{{expected ','}} + // expected-error@+4{{expected pointer in 'attach' clause, type is 'char'}} + // expected-error@+3{{OpenACC sub-array is not allowed here}} + // expected-note@+2{{expected variable of pointer type}} // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}} #pragma acc serial attach(s.array[s.value] s.array[s.value :5] ), seq for(;;){} - // expected-warning@+2{{OpenACC clause 'attach' not yet implemented, clause ignored}} // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}} -#pragma acc serial attach(s.array[s.value : 5], s.value), seq +#pragma acc serial attach(IsPointer), seq for(;;){} // expected-error@+3{{expected ','}} diff --git a/clang/test/SemaOpenACC/compute-construct-attach-clause.c b/clang/test/SemaOpenACC/compute-construct-attach-clause.c new file mode 100644 index 00000000000000..de735308528adb --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-attach-clause.c @@ -0,0 +1,61 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct S { + int IntMem; + int *PtrMem; +}; + +void uses() { + int LocalInt; + int *LocalPtr; + int Array[5]; + int *PtrArray[5]; + struct S s; + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(LocalInt) + while (1); + + // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}} +#pragma acc parallel attach(&LocalInt) + while (1); + +#pragma acc serial attach(LocalPtr) + while (1); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int[5]'}} +#pragma acc kernels attach(Array) + while (1); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(Array[0]) + while (1); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel attach(Array[0:1]) + while (1); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int *[5]'}} +#pragma acc parallel attach(PtrArray) + while (1); + +#pragma acc parallel attach(PtrArray[0]) + while (1); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel attach(PtrArray[0:1]) + while (1); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'struct S'}} +#pragma acc parallel attach(s) + while (1); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(s.IntMem) + while (1); + +#pragma acc parallel attach(s.PtrMem) + while (1); +} diff --git a/clang/test/SemaOpenACC/compute-construct-attach-clause.cpp b/clang/test/SemaOpenACC/compute-construct-attach-clause.cpp new file mode 100644 index 00000000000000..a89d346c2645ab --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-attach-clause.cpp @@ -0,0 +1,120 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct S { + int IntMem; + int *PtrMem; + operator int*(); +}; + +void uses() { + int LocalInt; + int *LocalPtr; + int Array[5]; + int *PtrArray[5]; + struct S s; + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(LocalInt) + while (true); + +#pragma acc parallel attach(LocalPtr) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int[5]'}} +#pragma acc parallel attach(Array) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(Array[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel attach(Array[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int *[5]'}} +#pragma acc parallel attach(PtrArray) + while (true); + +#pragma acc parallel attach(PtrArray[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel attach(PtrArray[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'struct S'}} +#pragma acc parallel attach(s) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(s.IntMem) + while (true); + +#pragma acc parallel attach(s.PtrMem) + while (true); +} + +template<typename T, typename TPtr, typename TStruct, auto &R1> +void Templ() { + T SomeInt; + TPtr SomePtr; + T SomeIntArray[5]; + TPtr SomeIntPtrArray[5]; + TStruct SomeStruct; + + // expected-error@+2{{expected pointer in 'attach' clause, type is 'int'}} + // expected-note@#INST{{in instantiation of function template specialization}} +#pragma acc parallel attach(SomeInt) + while (true); + +#pragma acc parallel attach(SomePtr) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int[5]'}} +#pragma acc parallel attach(SomeIntArray) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(SomeIntArray[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel attach(SomeIntArray[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int *[5]'}} +#pragma acc parallel attach(SomeIntPtrArray) + while (true); + +#pragma acc parallel attach(SomeIntPtrArray[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel attach(SomeIntPtrArray[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'S'}} +#pragma acc parallel attach(SomeStruct) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(SomeStruct.IntMem) + while (true); + +#pragma acc parallel attach(SomeStruct.PtrMem) + while (true); + + // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}} +#pragma acc parallel attach(R1) + while (true); +} + +void inst() { + static constexpr int CEVar = 1; + Templ<int, int*, S, CEVar>(); // #INST +} diff --git a/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c new file mode 100644 index 00000000000000..e5d328eb0b28bc --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c @@ -0,0 +1,61 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct S { + int IntMem; + int *PtrMem; +}; + +void uses() { + int LocalInt; + int *LocalPtr; + int Array[5]; + int *PtrArray[5]; + struct S s; + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(LocalInt) + while (1); + + // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}} +#pragma acc parallel deviceptr(&LocalInt) + while (1); + +#pragma acc serial deviceptr(LocalPtr) + while (1); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}} +#pragma acc kernels deviceptr(Array) + while (1); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(Array[0]) + while (1); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel deviceptr(Array[0:1]) + while (1); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}} +#pragma acc parallel deviceptr(PtrArray) + while (1); + +#pragma acc parallel deviceptr(PtrArray[0]) + while (1); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel deviceptr(PtrArray[0:1]) + while (1); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'struct S'}} +#pragma acc parallel deviceptr(s) + while (1); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(s.IntMem) + while (1); + +#pragma acc parallel deviceptr(s.PtrMem) + while (1); +} diff --git a/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp new file mode 100644 index 00000000000000..83409c91d4818f --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp @@ -0,0 +1,120 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct S { + int IntMem; + int *PtrMem; + operator int*(); +}; + +void uses() { + int LocalInt; + int *LocalPtr; + int Array[5]; + int *PtrArray[5]; + struct S s; + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(LocalInt) + while (true); + +#pragma acc parallel deviceptr(LocalPtr) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}} +#pragma acc parallel deviceptr(Array) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(Array[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel deviceptr(Array[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}} +#pragma acc parallel deviceptr(PtrArray) + while (true); + +#pragma acc parallel deviceptr(PtrArray[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel deviceptr(PtrArray[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'struct S'}} +#pragma acc parallel deviceptr(s) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(s.IntMem) + while (true); + +#pragma acc parallel deviceptr(s.PtrMem) + while (true); +} + +template<typename T, typename TPtr, typename TStruct, auto &R1> +void Templ() { + T SomeInt; + TPtr SomePtr; + T SomeIntArray[5]; + TPtr SomeIntPtrArray[5]; + TStruct SomeStruct; + + // expected-error@+2{{expected pointer in 'deviceptr' clause, type is 'int'}} + // expected-note@#INST{{in instantiation of function template specialization}} +#pragma acc parallel deviceptr(SomeInt) + while (true); + +#pragma acc parallel deviceptr(SomePtr) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}} +#pragma acc parallel deviceptr(SomeIntArray) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(SomeIntArray[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel deviceptr(SomeIntArray[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}} +#pragma acc parallel deviceptr(SomeIntPtrArray) + while (true); + +#pragma acc parallel deviceptr(SomeIntPtrArray[0]) + while (true); + + // expected-error@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel deviceptr(SomeIntPtrArray[0:1]) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'S'}} +#pragma acc parallel deviceptr(SomeStruct) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(SomeStruct.IntMem) + while (true); + +#pragma acc parallel deviceptr(SomeStruct.PtrMem) + while (true); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel deviceptr(R1) + while (true); +} + +void inst() { + static constexpr int CEVar = 1; + Templ<int, int*, S, CEVar>(); // #INST +} diff --git a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp index d35f62adfe0794..e057678d924957 100644 --- a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp +++ b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp @@ -191,6 +191,17 @@ void NormalUses(float *PointerParam) { // CHECK-NEXT: WhileStmt // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: NullStmt + +#pragma acc parallel attach(PointerParam) deviceptr(PointerParam) + while (true); + // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel + // CHECK-NEXT: attach clause + // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *' + // CHECK-NEXT: deviceptr clause + // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: NullStmt } // This example is an error typically, but we want to make sure we're properly @@ -402,6 +413,17 @@ void TemplUses(T t, U u, T*PointerParam) { // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: NullStmt +#pragma acc parallel attach(PointerParam) deviceptr(PointerParam) + while (true); + // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel + // CHECK-NEXT: attach clause + // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *' + // CHECK-NEXT: deviceptr clause + // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: NullStmt + // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl{{.*}}EndMarker int EndMarker; @@ -604,6 +626,16 @@ void TemplUses(T t, U u, T*PointerParam) { // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: NullStmt +//#pragma acc parallel attach(PointerParam) deviceptr(PointerParam) + // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel + // CHECK-NEXT: attach clause + // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *' + // CHECK-NEXT: deviceptr clause + // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: NullStmt + // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl{{.*}}EndMarker } @@ -613,6 +645,8 @@ struct S { // CHECK: CXXRecordDecl{{.*}} implicit struct S int ThisMember; // CHECK-NEXT: FieldDecl{{.*}} ThisMember 'int' + int *ThisMemberPtr; + // CHECK-NEXT: FieldDecl{{.*}} ThisMemberPtr 'int *' int ThisMemberArray[5]; // CHECK-NEXT: FieldDecl{{.*}} ThisMemberArray 'int[5]' @@ -620,10 +654,11 @@ struct S { // CHECK-NEXT: CXXMethodDecl{{.*}} foo 'void ()' template<typename T> - void bar() { + void bar(T *PointerParam) { // CHECK-NEXT: FunctionTemplateDecl{{.*}}bar // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T - // CHECK-NEXT: CXXMethodDecl{{.*}} bar 'void ()' implicit-inline + // CHECK-NEXT: CXXMethodDecl{{.*}} bar 'void (T *)' implicit-inline + // CHECK-NEXT: ParmVarDecl{{.*}} PointerParam 'T *' // CHECK-NEXT: CompoundStmt #pragma acc parallel private(ThisMember, this->ThisMemberArray[1]) @@ -664,10 +699,28 @@ struct S { // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: NullStmt +#pragma acc parallel attach(PointerParam, this, this->ThisMemberPtr) deviceptr(PointerParam, this, ThisMemberPtr) + while (true); + // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel + // CHECK-NEXT: attach clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *' + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this + // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this + // CHECK-NEXT: deviceptr clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *' + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this + // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' implicit this + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: NullStmt + // Check Instantiations: - // CHECK-NEXT: CXXMethodDecl{{.*}} used bar 'void ()' implicit_instantiation implicit-inline + // CHECK-NEXT: CXXMethodDecl{{.*}} used bar 'void (int *)' implicit_instantiation implicit-inline // CHECK-NEXT: TemplateArgument type 'int' // CHECK-NEXT: BuiltinType{{.*}} 'int' + // CHECK-NEXT: ParmVarDecl{{.*}} PointerParam 'int *' // CHECK-NEXT: CompoundStmt // #pragma acc parallel private(ThisMember, this->ThisMemberArray[1]) @@ -704,6 +757,22 @@ struct S { // CHECK-NEXT: WhileStmt // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: NullStmt + +//#pragma acc parallel attach(PointerParam, this, this->ThisMemberPtr) deviceptr(PointerParam, this, ThisMemberPtr) + // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel + // CHECK-NEXT: attach clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *' + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this + // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this + // CHECK-NEXT: deviceptr clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *' + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this + // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' implicit this + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: NullStmt } }; @@ -906,7 +975,7 @@ void Inst() { TemplUses<CEVar, int, int[1]>({}, {}, &i); S s; - s.bar<int>(); + s.bar<int>(&i); STempl<int> stempl; stempl.bar<int>(); } diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 487f5785957685..6c07c4d2e30738 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2840,6 +2840,13 @@ void OpenACCClauseEnqueue::VisitCopyOutClause(const OpenACCCopyOutClause &C) { void OpenACCClauseEnqueue::VisitCreateClause(const OpenACCCreateClause &C) { VisitVarList(C); } +void OpenACCClauseEnqueue::VisitAttachClause(const OpenACCAttachClause &C) { + VisitVarList(C); +} +void OpenACCClauseEnqueue::VisitDevicePtrClause( + const OpenACCDevicePtrClause &C) { + VisitVarList(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
