diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index a4cde9ae8a665..8b6d3221aa066 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -196,6 +196,16 @@ class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr { Expr *IntExpr, SourceLocation EndLoc); }; +class OpenACCVectorLengthClause : public OpenACCClauseWithSingleIntExpr { + OpenACCVectorLengthClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + Expr *IntExpr, SourceLocation EndLoc); + +public: + static OpenACCVectorLengthClause * + Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, + Expr *IntExpr, SourceLocation EndLoc); +}; + template class OpenACCClauseVisitor { Impl &getDerived() { return static_cast(*this); } diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def index d1a95cbe61394..520e068f4ffd4 100644 --- a/clang/include/clang/Basic/OpenACCClauses.def +++ b/clang/include/clang/Basic/OpenACCClauses.def @@ -19,5 +19,6 @@ VISIT_CLAUSE(Default) VISIT_CLAUSE(If) VISIT_CLAUSE(Self) VISIT_CLAUSE(NumWorkers) +VISIT_CLAUSE(VectorLength) #undef VISIT_CLAUSE diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index b995edebe8eed..023722049732a 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -93,13 +93,15 @@ class SemaOpenACC : public SemaBase { } unsigned getNumIntExprs() const { - assert(ClauseKind == OpenACCClauseKind::NumWorkers && + assert((ClauseKind == OpenACCClauseKind::NumWorkers || + ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); return std::get(Details).IntExprs.size(); } ArrayRef getIntExprs() { - assert(ClauseKind == OpenACCClauseKind::NumWorkers && + assert((ClauseKind == OpenACCClauseKind::NumWorkers || + ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); return std::get(Details).IntExprs; } @@ -132,7 +134,8 @@ class SemaOpenACC : public SemaBase { } void setIntExprDetails(ArrayRef IntExprs) { - assert(ClauseKind == OpenACCClauseKind::NumWorkers && + assert((ClauseKind == OpenACCClauseKind::NumWorkers || + ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}}; } diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index 3fff02fa33f28..75334223e073c 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -103,6 +103,27 @@ OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc, OpenACCNumWorkersClause(BeginLoc, LParenLoc, IntExpr, EndLoc); } +OpenACCVectorLengthClause::OpenACCVectorLengthClause(SourceLocation BeginLoc, + SourceLocation LParenLoc, + Expr *IntExpr, + SourceLocation EndLoc) + : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::VectorLength, BeginLoc, + LParenLoc, IntExpr, EndLoc) { + assert((!IntExpr || IntExpr->isInstantiationDependent() || + IntExpr->getType()->isIntegerType()) && + "Condition expression type not scalar/dependent"); +} + +OpenACCVectorLengthClause * +OpenACCVectorLengthClause::Create(const ASTContext &C, SourceLocation BeginLoc, + SourceLocation LParenLoc, Expr *IntExpr, + SourceLocation EndLoc) { + void *Mem = C.Allocate(sizeof(OpenACCVectorLengthClause), + alignof(OpenACCVectorLengthClause)); + return new (Mem) + OpenACCVectorLengthClause(BeginLoc, LParenLoc, IntExpr, EndLoc); +} + //===----------------------------------------------------------------------===// // OpenACC clauses printing methods //===----------------------------------------------------------------------===// @@ -124,3 +145,8 @@ void OpenACCClausePrinter::VisitNumWorkersClause( const OpenACCNumWorkersClause &C) { OS << "num_workers(" << C.getIntExpr() << ")"; } + +void OpenACCClausePrinter::VisitVectorLengthClause( + const OpenACCVectorLengthClause &C) { + OS << "vector_length(" << C.getIntExpr() << ")"; +} diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index cc230909bd384..8138ae3a244a8 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2503,6 +2503,12 @@ void OpenACCClauseProfiler::VisitNumWorkersClause( Profiler.VisitStmt(Clause.getIntExpr()); } +void OpenACCClauseProfiler::VisitVectorLengthClause( + const OpenACCVectorLengthClause &Clause) { + assert(Clause.hasIntExpr() && + "vector_length clause requires a valid int expr"); + Profiler.VisitStmt(Clause.getIntExpr()); +} } // namespace void StmtProfiler::VisitOpenACCComputeConstruct( diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 9d1b73cb7a078..e5a8b285715b3 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -400,6 +400,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { case OpenACCClauseKind::If: case OpenACCClauseKind::Self: case OpenACCClauseKind::NumWorkers: + case OpenACCClauseKind::VectorLength: // The condition expression will be printed as a part of the 'children', // but print 'clause' here so it is clear what is happening from the dump. OS << " clause"; diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 096e0863ed47c..757417f75c963 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -960,7 +960,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( // TODO OpenACC: as we implement the 'rest' of the above, this 'if' should // be removed leaving just the 'setIntExprDetails'. - if (ClauseKind == OpenACCClauseKind::NumWorkers) + if (ClauseKind == OpenACCClauseKind::NumWorkers || + ClauseKind == OpenACCClauseKind::VectorLength) ParsedClause.setIntExprDetails(IntExpr.get()); break; diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 39cde677ecc87..190739fa02e93 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -92,6 +92,7 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, return false; } case OpenACCClauseKind::NumWorkers: + case OpenACCClauseKind::VectorLength: switch (DirectiveKind) { case OpenACCDirectiveKind::Parallel: case OpenACCDirectiveKind::Kernels: @@ -248,6 +249,25 @@ SemaOpenACC::ActOnClause(ArrayRef ExistingClauses, getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getIntExprs()[0], Clause.getEndLoc()); } + case OpenACCClauseKind::VectorLength: { + // 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; + + // There is no prose in the standard that says duplicates aren't allowed, + // but this diagnostic is present in other compilers, as well as makes + // sense. + if (checkAlreadyHasClauseOfKind(*this, ExistingClauses, Clause)) + return nullptr; + + assert(Clause.getIntExprs().size() == 1 && + "Invalid number of expressions for VectorLength"); + return OpenACCVectorLengthClause::Create( + getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(), + Clause.getIntExprs()[0], Clause.getEndLoc()); + } default: break; } diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 34b52a346a10d..ade33ec65038f 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11181,6 +11181,29 @@ void OpenACCClauseTransform::VisitNumWorkersClause( ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0], ParsedClause.getEndLoc()); } + +template +void OpenACCClauseTransform::VisitVectorLengthClause( + const OpenACCVectorLengthClause &C) { + Expr *IntExpr = const_cast(C.getIntExpr()); + assert(IntExpr && "vector_length clause constructed with invalid int expr"); + + ExprResult Res = Self.TransformExpr(IntExpr); + if (!Res.isUsable()) + return; + + Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid, + C.getClauseKind(), + C.getBeginLoc(), Res.get()); + if (!Res.isUsable()) + return; + + ParsedClause.setIntExprDetails(Res.get()); + NewClause = OpenACCVectorLengthClause::Create( + Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0], + ParsedClause.getEndLoc()); +} } // namespace template OpenACCClause *TreeTransform::TransformOpenACCClause( diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 6fff829cefb84..44e23919ea18e 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -11792,6 +11792,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { return OpenACCNumWorkersClause::Create(getContext(), BeginLoc, LParenLoc, IntExpr, EndLoc); } + case OpenACCClauseKind::VectorLength: { + SourceLocation LParenLoc = readSourceLocation(); + Expr *IntExpr = readSubExpr(); + return OpenACCVectorLengthClause::Create(getContext(), BeginLoc, LParenLoc, + IntExpr, EndLoc); + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: case OpenACCClauseKind::Seq: @@ -11820,7 +11826,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { case OpenACCClauseKind::Reduction: case OpenACCClauseKind::Collapse: case OpenACCClauseKind::Bind: - case OpenACCClauseKind::VectorLength: case OpenACCClauseKind::NumGangs: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 934b202ac3826..6dd87b5d200db 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -5014,7 +5014,7 @@ void ASTWriter::WriteSpecialDeclRecords(Sema &SemaRef) { if (!ModularCodegenDecls.empty()) Stream.EmitRecord(MODULAR_CODEGEN_DECLS, ModularCodegenDecls); - + // Write the record containing tentative definitions. RecordData TentativeDefinitions; AddLazyVectorEmiitedDecls(*this, SemaRef.TentativeDefinitions, @@ -5135,7 +5135,7 @@ void ASTWriter::WriteSpecialDeclRecords(Sema &SemaRef) { } if (!UndefinedButUsed.empty()) Stream.EmitRecord(UNDEFINED_BUT_USED, UndefinedButUsed); - + // Write all delete-expressions that we would like to // analyze later in AST. RecordData DeleteExprsToAnalyze; @@ -7663,6 +7663,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { AddStmt(const_cast(NWC->getIntExpr())); return; } + case OpenACCClauseKind::VectorLength: { + const auto *NWC = cast(C); + writeSourceLocation(NWC->getLParenLoc()); + AddStmt(const_cast(NWC->getIntExpr())); + return; + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: case OpenACCClauseKind::Seq: @@ -7691,7 +7697,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { case OpenACCClauseKind::Reduction: case OpenACCClauseKind::Collapse: case OpenACCClauseKind::Bind: - case OpenACCClauseKind::VectorLength: case OpenACCClauseKind::NumGangs: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index bd5f459eae074..ddf40f71701ed 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -893,11 +893,9 @@ void IntExprParsing() { #pragma acc parallel vector_length(5, 4) {} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}} #pragma acc parallel vector_length(5) {} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}} #pragma acc parallel vector_length(returns_int()) {} diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp index 09a90726c6939..8c1d643747996 100644 --- a/clang/test/ParserOpenACC/parse-clauses.cpp +++ b/clang/test/ParserOpenACC/parse-clauses.cpp @@ -12,11 +12,9 @@ void templ() { #pragma acc loop collapse(T::value) for(;;){} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}} #pragma acc parallel vector_length(T::value) for(;;){} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}} #pragma acc parallel vector_length(I) for(;;){} diff --git a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp index e3664460a6e5e..889025c26818d 100644 --- a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp +++ b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp @@ -77,6 +77,17 @@ void NormalUses() { // CHECK-NEXT: WhileStmt // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: CompoundStmt + +#pragma acc kernels vector_length(some_short()) + while(true){} + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: CallExpr{{.*}}'short' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' + // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt } template @@ -157,6 +168,25 @@ void TemplUses(T t, U u) { // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: CompoundStmt +#pragma acc kernels vector_length(u) + while(true){} + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + +#pragma acc parallel vector_length(U::value) + while(true){} + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + // Check the instantiated versions of the above. // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation // CHECK-NEXT: TemplateArgument type 'CorrectConvert' @@ -239,6 +269,25 @@ void TemplUses(T t, U u) { // CHECK-NEXT: WhileStmt // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '' .operator char + // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' + // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt } struct HasInt { diff --git a/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c new file mode 100644 index 0000000000000..cd85bdefb602d --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c @@ -0,0 +1,33 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +short getS(); + +void Test() { +#pragma acc parallel vector_length(1) + while(1); +#pragma acc kernels vector_length(1) + while(1); + + // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'serial' directive}} +#pragma acc serial vector_length(1) + while(1); + + struct NotConvertible{} NC; + // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type ('struct NotConvertible' invalid)}} +#pragma acc parallel vector_length(NC) + while(1); + +#pragma acc kernels vector_length(getS()) + while(1); + + struct Incomplete *SomeIncomplete; + + // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type ('struct Incomplete' invalid)}} +#pragma acc kernels vector_length(*SomeIncomplete) + while(1); + + enum E{A} SomeE; + +#pragma acc kernels vector_length(SomeE) + while(1); +} diff --git a/clang/test/SemaOpenACC/compute-construct-vector_length-clause.cpp b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.cpp new file mode 100644 index 0000000000000..f6c5dde1a0235 --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.cpp @@ -0,0 +1,133 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct NotConvertible{} NC; +struct Incomplete *SomeIncomplete; // #INCOMPLETE +enum E{} SomeE; +enum class E2{} SomeE2; + +struct CorrectConvert { + operator int(); +} Convert; + +struct ExplicitConvertOnly { + explicit operator int() const; // #EXPL_CONV +} Explicit; + +struct AmbiguousConvert{ + operator int(); // #AMBIG_INT + operator short(); // #AMBIG_SHORT + operator float(); +} Ambiguous; + +void Test() { +#pragma acc parallel vector_length(1) + while(1); +#pragma acc kernels vector_length(1) + while(1); + + // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type ('struct NotConvertible' invalid}} +#pragma acc parallel vector_length(NC) + while(1); + + // expected-error@+2{{OpenACC integer expression has incomplete class type 'struct Incomplete'}} + // expected-note@#INCOMPLETE{{forward declaration of 'Incomplete'}} +#pragma acc kernels vector_length(*SomeIncomplete) + while(1); + +#pragma acc parallel vector_length(SomeE) + while(1); + + // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type ('enum E2' invalid}} +#pragma acc kernels vector_length(SomeE2) + while(1); + +#pragma acc parallel vector_length(Convert) + while(1); + + // expected-error@+2{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}} + // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} +#pragma acc kernels vector_length(Explicit) + while(1); + + // expected-error@+3{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}} + // expected-note@#AMBIG_INT{{conversion to integral type 'int'}} + // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}} +#pragma acc parallel vector_length(Ambiguous) + while(1); +} + +struct HasInt { + using IntTy = int; + using ShortTy = short; + static constexpr int value = 1; + static constexpr AmbiguousConvert ACValue; + static constexpr ExplicitConvertOnly EXValue; + + operator char(); +}; + +template +void TestInst() { + + // expected-error@+1{{no member named 'Invalid' in 'HasInt'}} +#pragma acc parallel vector_length(HasInt::Invalid) + while (1); + + // expected-error@+2{{no member named 'Invalid' in 'HasInt'}} + // expected-note@#INST{{in instantiation of function template specialization 'TestInst' requested here}} +#pragma acc kernels vector_length(T::Invalid) + while (1); + + // expected-error@+3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}} + // expected-note@#AMBIG_INT{{conversion to integral type 'int'}} + // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}} +#pragma acc parallel vector_length(HasInt::ACValue) + while (1); + + // expected-error@+3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}} + // expected-note@#AMBIG_INT{{conversion to integral type 'int'}} + // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}} +#pragma acc kernels vector_length(T::ACValue) + while (1); + + // expected-error@+2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}} + // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} +#pragma acc parallel vector_length(HasInt::EXValue) + while (1); + + // expected-error@+2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}} + // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} +#pragma acc kernels vector_length(T::EXValue) + while (1); + +#pragma acc parallel vector_length(HasInt::value) + while (1); + +#pragma acc kernels vector_length(T::value) + while (1); + +#pragma acc parallel vector_length(HasInt::IntTy{}) + while (1); + +#pragma acc kernels vector_length(typename T::ShortTy{}) + while (1); + +#pragma acc parallel vector_length(HasInt::IntTy{}) + while (1); + +#pragma acc kernels vector_length(typename T::ShortTy{}) + while (1); + + HasInt HI{}; + T MyT{}; + +#pragma acc parallel vector_length(HI) + while (1); + +#pragma acc kernels vector_length(MyT) + while (1); +} + +void Inst() { + TestInst(); // #INST +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 011bec32bab31..cbc1d85bb33df 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2799,6 +2799,10 @@ void OpenACCClauseEnqueue::VisitNumWorkersClause( const OpenACCNumWorkersClause &C) { Visitor.AddStmt(C.getIntExpr()); } +void OpenACCClauseEnqueue::VisitVectorLengthClause( + const OpenACCVectorLengthClause &C) { + Visitor.AddStmt(C.getIntExpr()); +} } // namespace void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {