Index: include/clang/Basic/Attr.td =================================================================== --- include/clang/Basic/Attr.td +++ include/clang/Basic/Attr.td @@ -3199,6 +3199,32 @@ }]; } +def OMPDeclareVariantDecl : Attr { + let Spellings = [Pragma<"omp", "declare variant">]; + let Subjects = SubjectList<[Function]>; + let SemaHandler = 0; + let HasCustomParsing = 1; + let Documentation = [Undocumented]; + let Args = [ + StringArgument<"VariantFuncId"> + ]; + // 'VecVarNamedDecl' is the declaration of the vector variant, discovered + // during semantic analysis. It is required during code generation (to verify + // the signature). + let AdditionalMembers = [{ + void printPrettyPragma(raw_ostream & OS, const PrintingPolicy &Policy) + const { + OS << "(" << getVariantFuncId().str() << ")"; + } + private: + NamedDecl *VecVarNamedDecl; + + public: + void setVecVarNamedDecl(NamedDecl *in) {VecVarNamedDecl = in;} + NamedDecl *getVecVarNamedDecl() const {return VecVarNamedDecl;} + }]; +} + def OMPDeclareTargetDecl : InheritableAttr { let Spellings = [Pragma<"omp", "declare target">]; let SemaHandler = 0; Index: include/clang/Basic/DiagnosticParseKinds.td =================================================================== --- include/clang/Basic/DiagnosticParseKinds.td +++ include/clang/Basic/DiagnosticParseKinds.td @@ -1176,8 +1176,8 @@ "expected identifier specifying the name of the 'omp critical' directive">; def err_omp_expected_reduction_identifier : Error< "expected identifier or one of the following operators: '+', '-', '*', '&', '|', '^', '&&', or '||'">; -def err_omp_decl_in_declare_simd : Error< - "function declaration is expected after 'declare simd' directive">; +def err_omp_decl_after_declare_directive : Error< + "function declaration is expected after 'declare %select{simd|variant}0' directive">; def err_omp_unknown_map_type : Error< "incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">; def err_omp_unknown_map_type_modifier : Error< @@ -1201,8 +1201,12 @@ def warn_omp_more_one_device_type_clause : Warning< "more than one 'device_type' clause is specified">, InGroup; +def err_omp_expected_declared_func_ident : Error< + "no declaration for '%0', only declared functions can be used as ">; +def err_omp_expected_unqualified_declare_variant_id : Error< + "expected unqualified-id representing 'vector-variant-id'">; -// Pragma loop support. +// Pragmr loop support. def err_pragma_loop_missing_argument : Error< "missing argument; expected %select{an integer value|" "'enable'%select{|, 'full'}1%select{|, 'assume_safety'}2 or 'disable'}0">; Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -9172,10 +9172,10 @@ "the 'copyprivate' clause must not be used with the 'nowait' clause">; def note_omp_nowait_clause_here : Note< "'nowait' clause is here">; -def err_omp_single_decl_in_declare_simd : Error< - "single declaration is expected after 'declare simd' directive">; +def err_omp_single_decl_after_declare_directive : Error< + "single declaration is expected after 'declare %select{simd|variant}0' directive">; def err_omp_function_expected : Error< - "'#pragma omp declare simd' can only be applied to functions">; + "'#pragma omp declare %select{simd|variant}0' can only be applied to functions">; def err_omp_wrong_cancel_region : Error< "one of 'for', 'parallel', 'sections' or 'taskgroup' is expected">; def err_omp_parent_cancel_region_nowait : Error< @@ -9367,6 +9367,8 @@ def warn_omp_declare_target_after_first_use : Warning< "declaration marked as declare target after first use, it may lead to incorrect results">, InGroup; +def err_omp_declare_variant_template_decl : Error< + "'#pragma omp declare variant' can only be used to decorate instantiated templates">; } // end of OpenMP category let CategoryName = "Related Result Type Issue" in { Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -236,6 +236,7 @@ OPENMP_DIRECTIVE(distribute) OPENMP_DIRECTIVE_EXT(declare_target, "declare target") OPENMP_DIRECTIVE_EXT(end_declare_target, "end declare target") +OPENMP_DIRECTIVE_EXT(declare_variant, "declare variant") OPENMP_DIRECTIVE_EXT(distribute_parallel_for, "distribute parallel for") OPENMP_DIRECTIVE_EXT(distribute_parallel_for_simd, "distribute parallel for simd") OPENMP_DIRECTIVE_EXT(distribute_simd, "distribute simd") Index: include/clang/Parse/Parser.h =================================================================== --- include/clang/Parse/Parser.h +++ include/clang/Parse/Parser.h @@ -2834,6 +2834,11 @@ DeclGroupPtrTy ParseOMPDeclareSimdClauses(DeclGroupPtrTy Ptr, CachedTokens &Toks, SourceLocation Loc); + /// Parse clauses for '#pragma omp declare variant'. + DeclGroupPtrTy ParseOMPDeclareVariantClauses(DeclGroupPtrTy Ptr, + CachedTokens &Toks, + SourceLocation Loc); + bool parseDeclareVariantClauses(IdentifierLoc &VectorVariantId); /// Parse clauses for '#pragma omp declare target'. DeclGroupPtrTy ParseOMPDeclareTargetClauses(); /// Parse '#pragma omp end declare target'. Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -9477,6 +9477,11 @@ ArrayRef Alignments, ArrayRef Linears, ArrayRef LinModifiers, ArrayRef Steps, SourceRange SR); + /// Called on well-formed '\#pragma omp declare variant' after parsing of + /// the associated method/function. + DeclGroupPtrTy ActOnOpenMPDeclareVariantDirective( + DeclGroupPtrTy DG, IdentifierLoc &VectorVariantId, SourceRange SR); + OMPClause *ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr, SourceLocation StartLoc, Index: lib/Basic/OpenMPKinds.cpp =================================================================== --- lib/Basic/OpenMPKinds.cpp +++ lib/Basic/OpenMPKinds.cpp @@ -838,6 +838,7 @@ break; } break; + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_unknown: @@ -1075,6 +1076,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_requires: Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -9545,6 +9545,7 @@ CGM, ParentName, cast(E)); break; + case OMPD_declare_variant: case OMPD_parallel: case OMPD_for: case OMPD_parallel_for: @@ -10170,6 +10171,7 @@ RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait : OMPRTL__tgt_target_data_update; break; + case OMPD_declare_variant: case OMPD_parallel: case OMPD_for: case OMPD_parallel_for: Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -15,6 +15,7 @@ #include "clang/Parse/ParseDiagnostic.h" #include "clang/Parse/Parser.h" #include "clang/Parse/RAIIObjectsForParser.h" +#include "clang/Sema/Lookup.h" #include "clang/Sema/Scope.h" #include "llvm/ADT/PointerIntPair.h" @@ -38,6 +39,7 @@ OMPD_target_enter, OMPD_target_exit, OMPD_update, + OMPD_variant, OMPD_distribute_parallel, OMPD_teams_distribute_parallel, OMPD_target_teams_distribute_parallel, @@ -80,6 +82,7 @@ .Case("reduction", OMPD_reduction) .Case("update", OMPD_update) .Case("mapper", OMPD_mapper) + .Case("variant", OMPD_variant) .Default(OMPD_unknown); } @@ -92,6 +95,7 @@ {OMPD_declare, OMPD_reduction, OMPD_declare_reduction}, {OMPD_declare, OMPD_mapper, OMPD_declare_mapper}, {OMPD_declare, OMPD_simd, OMPD_declare_simd}, + {OMPD_declare, OMPD_variant, OMPD_declare_variant}, {OMPD_declare, OMPD_target, OMPD_declare_target}, {OMPD_distribute, OMPD_parallel, OMPD_distribute_parallel}, {OMPD_distribute_parallel, OMPD_for, OMPD_distribute_parallel_for}, @@ -743,6 +747,51 @@ return IsError; } +/// Parses clauses for 'declare variant' directive. +/// (vec-var-id) match(context-selector-specification) +/// context-selector-specification: +/// TODO Add support for matching context selector +/// +/// Returns 'true' if there were syntax errors, 'false' otherwise. +/// Note that similar method for '#pragma omp simd' is _not_ a member of the +/// Parser - it's a static function local to this file instead. However, this +/// method accesses the internals of the Parser, so it makes more sense to +/// make it a member. +bool Parser::parseDeclareVariantClauses(IdentifierLoc &VectorVariantId) { + const Token &Tok = getCurToken(); + + // Parse (variant-func-id) + BalancedDelimiterTracker VecVarLeftParen(*this, tok::l_paren); + if (!Tok.is(tok::l_paren)) { + Diag(Tok, diag::err_expected_lparen_after) << "#pragma omp declare variant"; + return true; + } + VecVarLeftParen.consumeOpen(); + + if (Tok.isNot(tok::identifier)) { + Diag(Tok, diag::err_omp_expected_unqualified_declare_variant_id); + SkipUntil(tok::r_paren, StopAtSemi); + return true; + } + VectorVariantId = *ParseIdentifierLoc(); + + // Verify that 'variant-func-id' has already been declared + auto Identinfo = PP.getIdentifierInfo(VectorVariantId.Ident->getName()); + LookupResult Lookup(Actions, Identinfo, VectorVariantId.Loc, + Sema::LookupOrdinaryName); + CXXScopeSpec SS; + if (!Actions.LookupParsedName(Lookup, getCurScope(), &SS)) { + Diag(VectorVariantId.Loc, diag::err_omp_expected_declared_func_ident) + << VectorVariantId.Ident->getName(); + return true; + } + + if (VecVarLeftParen.consumeClose()) + return true; + + return false; +} + /// Parse clauses for '#pragma omp declare simd'. Parser::DeclGroupPtrTy Parser::ParseOMPDeclareSimdClauses(Parser::DeclGroupPtrTy Ptr, @@ -782,6 +831,34 @@ LinModifiers, Steps, SourceRange(Loc, EndLoc)); } +/// Parse clauses for '#pragma omp declare variant' +Parser::DeclGroupPtrTy +Parser::ParseOMPDeclareVariantClauses(Parser::DeclGroupPtrTy Ptr, + CachedTokens &Toks, SourceLocation Loc) { + PP.EnterToken(Tok, /*IsReinject*/ true); + PP.EnterTokenStream(Toks, /*DisableMacroExpansion=*/true, + /*IsReinject*/ true); + // Consume the previously pushed token. + ConsumeAnyToken(/*ConsumeCodeCompletionTok=*/true); + + FNContextRAII FnContext(*this, Ptr); + IdentifierLoc VectorVariantId; + bool IsError = parseDeclareVariantClauses(VectorVariantId); + // Need to check for extra tokens. + if (Tok.isNot(tok::annot_pragma_openmp_end)) { + Diag(Tok, diag::warn_omp_extra_tokens_at_eol) + << getOpenMPDirectiveName(OMPD_declare_variant); + while (Tok.isNot(tok::annot_pragma_openmp_end)) + ConsumeAnyToken(); + } + // Skip the last annot_pragma_openmp_end. + SourceLocation EndLoc = ConsumeAnnotationToken(); + if (IsError) + return Ptr; + return Actions.ActOnOpenMPDeclareVariantDirective(Ptr, VectorVariantId, + SourceRange(Loc, EndLoc)); +} + /// Parsing of simple OpenMP clauses like 'default' or 'proc_bind'. /// /// default-clause: @@ -1103,7 +1180,16 @@ } break; } - case OMPD_declare_simd: { + case OMPD_declare_simd: + LLVM_FALLTHROUGH; + case OMPD_declare_variant: { + // The syntax is: + // { #pragma omp declare simd } + // + // and: + // { #pragma omp declare variant } + // + // // The syntax is: // { #pragma omp declare simd } // @@ -1133,10 +1219,13 @@ } } if (!Ptr) { - Diag(Loc, diag::err_omp_decl_in_declare_simd); + unsigned directive = (OMPD_declare_simd == DKind) ? 0 : 1; + Diag(Loc, diag::err_omp_decl_after_declare_directive) << directive; return DeclGroupPtrTy(); } - return ParseOMPDeclareSimdClauses(Ptr, Toks, Loc); + return (OMPD_declare_simd == DKind) + ? ParseOMPDeclareSimdClauses(Ptr, Toks, Loc) + : ParseOMPDeclareVariantClauses(Ptr, Toks, Loc); } case OMPD_declare_target: { SourceLocation DTLoc = ConsumeAnyToken(); @@ -1569,6 +1658,7 @@ break; } case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_requires: Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -3438,6 +3438,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_requires: @@ -4506,6 +4507,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_requires: llvm_unreachable("OpenMP Directive is not allowed"); case OMPD_unknown: @@ -4645,7 +4647,8 @@ return DeclGroupPtrTy(); if (!DG.get().isSingleDecl()) { - Diag(SR.getBegin(), diag::err_omp_single_decl_in_declare_simd); + Diag(SR.getBegin(), diag::err_omp_single_decl_after_declare_directive) + << /*declare simd*/ 0; return DG; } Decl *ADecl = DG.get().getSingleDecl(); @@ -4654,7 +4657,8 @@ auto *FD = dyn_cast(ADecl); if (!FD) { - Diag(ADecl->getLocation(), diag::err_omp_function_expected); + Diag(ADecl->getLocation(), diag::err_omp_function_expected) + << /*declare simd*/ 0; return DeclGroupPtrTy(); } @@ -4879,6 +4883,54 @@ return ConvertDeclToDeclGroup(ADecl); } +Sema::DeclGroupPtrTy Sema::ActOnOpenMPDeclareVariantDirective( + DeclGroupPtrTy DG, IdentifierLoc &VectorVariantId, SourceRange SR) { + assert(!DG.get().isNull() && "Missing function declaration after pragma"); + + // '#pragma omp declare variant' is only allowed for template instantiations. + // It's meaning for template declarations would be ambigous. + Decl *D = *DG.get().begin(); + if (D->isTemplateDecl()) { + Diag(SR.getBegin(), diag::err_omp_declare_variant_template_decl); + return DG; + } + + if (!DG.get().isSingleDecl()) { + Diag(SR.getBegin(), diag::err_omp_single_decl_after_declare_directive) + << /*declare variant*/ 1; + return DG; + } + + Decl *ADecl = DG.get().getSingleDecl(); + if (auto *FTD = dyn_cast(ADecl)) + ADecl = FTD->getTemplatedDecl(); + + auto *FD = dyn_cast(ADecl); + if (!FD) { + Diag(ADecl->getLocation(), diag::err_omp_function_expected) + << /*declare variant*/ 1; + return DeclGroupPtrTy(); + } + + // Look up for the declaration of the variant. + StringRef FuncName = VectorVariantId.Ident->getName(); + DeclarationName VecVarDN(VectorVariantId.Ident); + NamedDecl *VecVarND = LookupSingleName(getCurScope(), VecVarDN, SR.getBegin(), + Sema::LookupOrdinaryName); + + // If the declaration fo the variant is found, add the + // attribute. Otherwise, raise an error. + if (VecVarND) { + auto *NewAttr = + OMPDeclareVariantDeclAttr::CreateImplicit(Context, FuncName, SR); + NewAttr->setVecVarNamedDecl(VecVarND); + ADecl->addAttr(NewAttr); + } else + Diag(VectorVariantId.Loc, diag::err_undeclared_var_use) << FuncName; + + return ConvertDeclToDeclGroup(ADecl); +} + StmtResult Sema::ActOnOpenMPParallelDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, @@ -9885,6 +9937,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_teams: @@ -9953,6 +10006,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_teams: @@ -10022,6 +10076,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_simd: @@ -10088,6 +10143,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_simd: @@ -10155,6 +10211,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_simd: @@ -10221,6 +10278,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_simd: @@ -10286,6 +10344,7 @@ case OMPD_declare_reduction: case OMPD_declare_mapper: case OMPD_declare_simd: + case OMPD_declare_variant: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_simd: Index: test/AST/dump.cpp =================================================================== --- test/AST/dump.cpp +++ test/AST/dump.cpp @@ -81,7 +81,7 @@ } #pragma omp end declare target -// CHECK: `-FunctionDecl {{.+}} line:[[@LINE-6]]:5 bar 'int ()' +// CHECK: |-FunctionDecl {{.+}} line:[[@LINE-6]]:5 bar 'int ()' // CHECK-NEXT: |-CompoundStmt {{.+}} // CHECK-NEXT: | |-DeclStmt {{.+}} // CHECK-NEXT: | | `-VarDecl {{.+}} col:7 used f 'int' @@ -89,3 +89,13 @@ // CHECK-NEXT: | `-ImplicitCastExpr {{.+}} 'int' // CHECK-NEXT: | `-DeclRefExpr {{.+}} 'int' lvalue Var {{.+}} 'f' 'int' // CHECK-NEXT: `-OMPDeclareTargetDeclAttr {{.+}} <> Implicit MT_To + +void vector_fez_v1(); +void vector_fez_v2(); +#pragma omp declare variant (vector_fez_v1) +#pragma omp declare variant (vector_fez_v2) +void fez(); + +// CHECK: `-FunctionDecl {{.+}} col:6 fez 'void ()' +// CHECK-NEXT: |-OMPDeclareVariantDeclAttr {{.+}} Implicit "vector_fez_v2" +// CHECK: `-OMPDeclareVariantDeclAttr {{.+}} Implicit "vector_fez_v1 Index: test/OpenMP/declare_variant_ast_print.c =================================================================== --- /dev/null +++ test/OpenMP/declare_variant_ast_print.c @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s +// RUN: %clang_cc1 -fopenmp -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s + +// RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s +// RUN: %clang_cc1 -fopenmp-simd -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s +// expected-no-diagnostics + +void vector_add_1(float *d, int s1, float *s2, double b[]); +void vector_add_2(float *d, int s1, float *s2, double b[]); +void vector_add_3(float *d, int s1, float *s2, double b[]); +void vector_add_4(float *d, int s1, float *s2, double b[]); + +#pragma omp declare variant (vector_add_1) +#pragma omp declare variant (vector_add_2) +#pragma omp declare variant (vector_add_3) +#pragma omp declare variant (vector_add_4) +void add(float *d, int s1, float *s2, double b[]); + +// CHECK: void vector_add_1(float *d, int s1, float *s2, double b[]); +// CHECK-NEXT: void vector_add_2(float *d, int s1, float *s2, double b[]); +// CHECK-NEXT: void vector_add_3(float *d, int s1, float *s2, double b[]); +// CHECK-NEXT: void vector_add_4(float *d, int s1, float *s2, double b[]); + +// CHECK: #pragma omp declare variant(vector_add_4) +// CHECK-NEXT: #pragma omp declare variant(vector_add_3) +// CHECK-NEXT: #pragma omp declare variant(vector_add_2) +// CHECK-NEXT: #pragma omp declare variant(vector_add_1) +// CHECK-NEXT: void add(float *d, int s1, float *s2, double b[]); Index: test/OpenMP/declare_variant_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/declare_variant_messages.cpp @@ -0,0 +1,84 @@ +// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c++ -std=c++11 -fms-extensions -Wno-pragma-pack %s + +// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp-simd -x c++ -std=c++11 -fms-extensions -Wno-pragma-pack %s + +// The dummy vector variant +void vector_foo(); + +//============================================================================= +// Basic errors +//============================================================================= +// This is (and should be) perfectly fine for the parser, although semantically +// not valid OpenMP 5. +#pragma omp declare variant(vector_foo) +void foo(); + +// expected-error@+1 {{expected an OpenMP directive}} +#pragma omp declare + +// expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}} +#pragma omp declare variant(vector_foo) +int a; +// expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}} +#pragma omp declare variant (vector_foo) +#pragma omp threadprivate(a) +int var; +#pragma omp threadprivate(var) + +// expected-error@+3 {{expected an OpenMP directive}} +// expected-error@+1 {{function declaration is expected after 'declare variant' directive}} +#pragma omp declare variant +#pragma omp declare + +// expected-error@+1 {{single declaration is expected after 'declare variant' directive}} +#pragma omp declare variant(vector_foo) +int b, c; + +//============================================================================= +// Syntax/semantic error: everything excluding linear/aligned/uniform +//============================================================================= +// expected-error@+1 {{expected an OpenMP directive}} +#pragma omp variant +// expected-error@+1 {{expected '(' after '#pragma omp declare variant'}} +#pragma omp declare variant +// expected-warning@+2 {{extra tokens at the end of '#pragma omp declare variant' are ignored}} +// expected-error@+1 {{expected '(' after '#pragma omp declare variant'}} +#pragma omp declare variant ) +// expected-error@+1 {{expected unqualified-id representing 'vector-variant-id'}} +#pragma omp declare variant() +// expected-warning@+2 {{extra tokens at the end of '#pragma omp declare variant' are ignored}} +// expected-error@+1 {{no declaration for 'undeclared_func', only declared functions can be used as }} +#pragma omp declare variant(undeclared_func) +// expected-error@+2 {{expected ')'}} +// expected-note@+1 {{to match this '('}} +#pragma omp declare variant(vector_foo +void foo(); + +//============================================================================= +// Templates +//============================================================================= +// expected-error@+1 {{'#pragma omp declare variant' can only be used to decorate instantiated templates}} +#pragma omp declare variant(vector_foo) +template +void h(C *hp, C *hp2, C *hq, C *lin) { + b = 0; +} + +#pragma omp declare variant(vector_foo) +template <> +void h(int *hp, int *hp2, int *hq, int *lin) { + h((float *)hp, (float *)hp2, (float *)hq, (float *)lin); +} + +//============================================================================= +// 'declare variant' embedded in namespaces +//============================================================================= +namespace N { + // expected-error@+1 {{function declaration is expected after 'declare variant' directive}} + #pragma omp declare variant +} +// expected-error@+1 {{function declaration is expected after 'declare variant' directive}} +#pragma omp declare variant +// Now verify that the parser recovered fine from the previous error and identifies another one correctly +// expected-error@+1 {{function declaration is expected after 'declare variant' directive}} +#pragma omp declare variant