[OpenMP] Support construct trait set for Clang

[cchen]

Summary

This patch supports construct trait set selector by using the existed
declare variant infrastructure inside OMPContext and simd selector is
currently not supported. The goal of this patch is to pass the declare variant
test inside sollve test suite.

[cchen]

clang/lib/AST/OpenMPClause.cpp

2344–2345

I'm now removing this line since the properties are already added at line 2334 and not removing this line would lead to an extra iteration in OMPContext.cpp (check out the second diff), and the code will always return false even if the construct trait properties are added.

Diff for OpenMPClause.cpp

2333       for (const OMPTraitProperty &Property : Selector.Properties)
2334         VMI.addTrait(Set.Kind, Property.Kind, Property.RawString, ScorePtr);
2335
2336       if (Set.Kind != TraitSet::construct)
2337         continue;
2338
2339       // TODO: This might not hold once we implement SIMD properly.
2340       assert(Selector.Properties.size() == 1 &&
2341              Selector.Properties.front().Kind ==
2342                  getOpenMPContextTraitPropertyForSelector(
2343                      Selector.Kind) &&
2344              "Ill-formed construct selector!");
2345
2346       // VMI.ConstructTraits.push_back(Selector.Properties.front().Kind);

Diff for lib/Frontend/OpenMP/OMPContext.cpp:

224     unsigned ConstructIdx = 0, NoConstructTraits = Ctx.ConstructTraits.size();
225     for (TraitProperty Property : VMI.ConstructTraits) {
226       assert(getOpenMPContextTraitSetForProperty(Property) ==
227                  TraitSet::construct &&
228              "Variant context is ill-formed!");
229
230       // Verify the nesting.
231       bool FoundInOrder = false;
232       while (!FoundInOrder && ConstructIdx != NoConstructTraits)
233         FoundInOrder = (Ctx.ConstructTraits[ConstructIdx++] == Property);
234       if (ConstructMatches)
235         ConstructMatches->push_back(ConstructIdx - 1);
236
237       Optional<bool> Result = HandleTrait(Property, FoundInOrder);
238       if (Result.hasValue())
239         return Result.getValue();
240
241       if (!FoundInOrder) {
242         LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] Construct property "
243                           << getOpenMPContextTraitPropertyName(Property, "")
244                           << " was not nested properly.\n");
245         return false;
246       }
247
248       // TODO: Verify SIMD
249     }

[cchen]

clang/test/OpenMP/declare_variant_construct_codegen_1.c

24

This test is from sollve test suite: https://github.com/SOLLVE/sollve_vv/blob/master/tests/5.0/declare_variant/test_declare_variant.c

[jdoerfert]

Cool!

We certainly need more tests but the design you came up with is certainly nicer than what I had in mind.
One thing I wanted to do though is keep track of the constructs in the OpenMPIRBuilder instead. So move the

/// Vector of declare variant construct traits.
SmallVector<llvm::omp::TraitProperty, 8> ConstructTraits;

into OpenMPIRBuilder and add/delete the things there. It's not a conceptual difference but makes it easier
to opt-in for Flang later.

I also left some comments below. I don't think there is anything major though.

clang/lib/AST/OpenMPClause.cpp

2344–2345

I'm not surprised some things need adjustment. Never tested construct trait matching much as we did not expose it via clang. Just delete the line if you think it's wrong. Also consider deleting the code before if it's not needed.

clang/lib/Sema/SemaOpenMP.cpp

737

Modification was made to make it clear what happens but now I think this is not needed.

See below first. Maybe the interface should be "handleConstructTrait(Property, bool Insert)" and if Insert is false you verify it's the last one and delete it. You'd need to traverse them in opposite order though to verify, unsure if verification is therefore strictly necessary.

3900–3902

4497

I know think you need to call something like addDeclareVariantConstructTrait again and do pop_back() instead of push_back() for each trait matching the DKind. Clearing, or erasing a fixed amount of traits, does not work.

6845

Maybe we should pass the traits to the constructor, makes it clearer they have to be passed.

[cchen]

One thing I wanted to do though is keep track of the constructs in the OpenMPIRBuilder instead. So move the

/// Vector of declare variant construct traits.
SmallVector<llvm::omp::TraitProperty, 8> ConstructTraits;

into OpenMPIRBuilder and add/delete the things there. It's not a conceptual difference but makes it easier
to opt-in for Flang later.

We don't have any OpenMPIRBuilder instance inside SemaOpenMP yet and the OpenMPIRBuilder constructor requires a Module parameter so I'm not sure how to achieve this, can you give me some suggestions?

clang/lib/Sema/SemaOpenMP.cpp

737

For the delete part, I'm not sure how passing Property would work since I'm not deleting all the construct trait properties inside ActOnOpenMPRegionEnd function.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

737

I think I've figured out a way to achieve what you suggest.

4497

So you are saying that instead of removing all the construct property inside ActOnOpenMPRegionEnd, we should delete those construct properties more fine-grained? I'd like to do so and has tried putting the insert/delete logic inside more specific function such as ActOnOpenMPParllell,

[cchen]

Rebase and fix based on suggestions

Haven't moved ConstructTrait to IRBuilder

[jdoerfert]

clang/lib/AST/OpenMPClause.cpp

2487

Nit: no braces.

clang/lib/Parse/ParseOpenMP.cpp

2052

{} might do it but this is fine too.

clang/lib/Sema/SemaOpenMP.cpp

746

So, SmalVector has a pop_back_val which we could use to verify the properties are the ones we expect. However, we would need to first put them into a vector and if insert is false we would need to reverse the vector.
Maybe that's the way to go:

1. Where you call handleConstructTrait right now, just put the traits in a vector.
2. Make it handleConstructTraits and take a vector of construct traits.
3. For insert == true, just append them all.
4. For insert == falser reverse the vector and do pop_back_val with an assertion that they match the result of vector.pop_back_val();

Does that make sense?

(Apologies for making you rewrite this twice.)

[cchen]

Fix based on feedback (wait for comment about moving ConstructTrait to IRBuilder)

[cchen]

Remove braces

[jdoerfert]

clang/lib/Sema/SemaOpenMP.cpp

746

This is not what I meant and I doubt this works. Let me try to explain with an example:

handleDeclareVariantConstructTrait for a DSKind parallel_for will first add a parallel then a for construct trait.
When we leave the parallel for we should pop them of our construct trait stack again, and preferably verify we do see the ones we expect.
So let's assume the stack was like this before

[ teams ]
// enter parallel for code generation
[ teams, parallel, for]
...
// exit parallel for code generation

at this point we want to pop for and then parallel from the end.

If the last one on the stack was not for or the second to last was not parallel something went wrong.
Since handleDeclareVariantConstructTrait does give use the traits "in order", thus [parallel, for], we need to do something different to verify this is the suffix of the stack.

I proposed to accumulate all traits in handleDeclareVariantConstructTrait rather than passing them to handleConstructTrait one by one.
So there would be a single call to handleConstructTrait at the end with all the traits, here [parallel, for].
If we are entering the scope we just append them to the stack: [teams] + [parallel, for] = [teams, parallel, for]
If we are exiting the scope (Insert == false, or name it ScopeEntry instead), we reverse the trait set and compare it against what we push from the stack:

stack = [teams, parallel, for] 
trait_set = [parallel, for] (passed to `handleConstructTrait` by `handleDeclareVariantConstructTrait) 
verification:
  for (trait in reverse(trait_set)) {
    stack_last = stack.pop_back_val();
    assert(trait == stack_last && "Something left a trait on the stack!");
  }

Does this make sense?

[cchen]

Fix based on suggestions

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

746

Thanks for your further explanation and patience, it really helps a lot for me to understand!

[jdoerfert]

Great! This should work. Though we need actual tests ;)

I also provided you with two minor edits below to make it easier to read, no functional change though.

clang/lib/Sema/SemaOpenMP.cpp

737–738

740

[cchen]

Add tests and fix coding style

[jdoerfert]

LGTM, thanks :)

[MaskRay]

clang/lib/Sema/SemaOpenMP.cpp

741

I fixed -Wunused-variable in c5f480fcbec04f46a9cfcad08914665ff83d8d8a

clang/include/clang/AST/OpenMPClause.h

public:
   /// Print a human readable representation into \p OS.
   void print(llvm::raw_ostream &OS, const PrintingPolicy &Policy) const;
 };
 llvm::raw_ostream &operator<<(llvm::raw_ostream &OS, const OMPTraitInfo &TI);
 llvm::raw_ostream &operator<<(llvm::raw_ostream &OS, const OMPTraitInfo *TI);
 
 /// Clang specific specialization of the OMPContext to lookup target features.
 struct TargetOMPContext final : public llvm::omp::OMPContext {
-
   TargetOMPContext(ASTContext &ASTCtx,
                    std::function<void(StringRef)> &&DiagUnknownTrait,
-                   const FunctionDecl *CurrentFunctionDecl);
+                   const FunctionDecl *CurrentFunctionDecl,
+                   ArrayRef<llvm::omp::TraitProperty> ConstructTraits);
+
   virtual ~TargetOMPContext() = default;
 
   /// See llvm::omp::OMPContext::matchesISATrait
   bool matchesISATrait(StringRef RawString) const override;
 
 private:
   std::function<bool(StringRef)> FeatureValidityCheck;
   std::function<void(StringRef)> DiagUnknownTrait;

clang/lib/AST/OpenMPClause.cpp

for (const OMPTraitSelector &Selector : Set.Selectors) {
       if (Set.Kind != TraitSet::construct)
         continue;
 
       // TODO: This might not hold once we implement SIMD properly.
       assert(Selector.Properties.size() == 1 &&
              Selector.Properties.front().Kind ==
                  getOpenMPContextTraitPropertyForSelector(
                      Selector.Kind) &&
              "Ill-formed construct selector!");
-
-      VMI.ConstructTraits.push_back(Selector.Properties.front().Kind);
     }
   }
 }
 
 void OMPTraitInfo::print(llvm::raw_ostream &OS,
                          const PrintingPolicy &Policy) const {
   bool FirstSet = true;
   for (const OMPTraitSet &Set : Sets) {
     if (!FirstSet)

 }
 llvm::raw_ostream &clang::operator<<(llvm::raw_ostream &OS,
                                      const OMPTraitInfo *TI) {
   return TI ? OS << *TI : OS;
 }
 
 TargetOMPContext::TargetOMPContext(
     ASTContext &ASTCtx, std::function<void(StringRef)> &&DiagUnknownTrait,
-    const FunctionDecl *CurrentFunctionDecl)
+    const FunctionDecl *CurrentFunctionDecl,
+    ArrayRef<llvm::omp::TraitProperty> ConstructTraits)
     : OMPContext(ASTCtx.getLangOpts().OpenMPIsDevice,
                  ASTCtx.getTargetInfo().getTriple()),
       FeatureValidityCheck([&](StringRef FeatureName) {
         return ASTCtx.getTargetInfo().isValidFeatureName(FeatureName);
       }),
       DiagUnknownTrait(std::move(DiagUnknownTrait)) {
   ASTCtx.getFunctionFeatureMap(FeatureMap, CurrentFunctionDecl);
+
+  for (llvm::omp::TraitProperty Property : ConstructTraits)
+    addTrait(Property);
 }
 
 bool TargetOMPContext::matchesISATrait(StringRef RawString) const {
   auto It = FeatureMap.find(RawString);
   if (It != FeatureMap.end())
     return It->second;
   if (!FeatureValidityCheck(RawString))
     DiagUnknownTrait(RawString);
   return false;
 }

clang/lib/Parse/ParseOpenMP.cpp

 
 static StringRef stringLiteralParser(Parser &P) {
   ExprResult Res = P.ParseStringLiteralExpression(true);
   return Res.isUsable() ? Res.getAs<StringLiteral>()->getString() : "";
 }
 
 static StringRef getNameFromIdOrString(Parser &P, Token &Tok,
                                        OMPContextLvl Lvl) {
-  if (Tok.is(tok::identifier)) {
+  if (Tok.is(tok::identifier) || Tok.is(tok::kw_for)) {
     llvm::SmallString<16> Buffer;
     StringRef Name = P.getPreprocessor().getSpelling(Tok, Buffer);
     (void)P.ConsumeToken();
     return Name;
   }
 
   if (tok::isStringLiteral(Tok.getKind()))
     return stringLiteralParser(P);
case OMPD_begin_declare_variant: {
     TI.getAsVariantMatchInfo(ASTCtx, VMI);
 
     std::function<void(StringRef)> DiagUnknownTrait = [this, Loc](
                                                           StringRef ISATrait) {
       // TODO Track the selector locations in a way that is accessible here to
       // improve the diagnostic location.
       Diag(Loc, diag::warn_unknown_begin_declare_variant_isa_trait) << ISATrait;
     };
-    TargetOMPContext OMPCtx(ASTCtx, std::move(DiagUnknownTrait),
-                            /* CurrentFunctionDecl */ nullptr);
+    TargetOMPContext OMPCtx(
+        ASTCtx, std::move(DiagUnknownTrait),
+        /* CurrentFunctionDecl */ nullptr,
+        /* ConstructTraits */ ArrayRef<llvm::omp::TraitProperty>());
 
     if (isVariantApplicableInContext(VMI, OMPCtx, /* DeviceSetOnly */ true)) {
       Actions.ActOnOpenMPBeginDeclareVariant(Loc, TI);
       break;
     }
 
     // Elide all the code till the matching end declare variant was found.
     unsigned Nesting = 1;

clang/lib/Sema/SemaOpenMP.cpp

private:
   /// omp_alloctrait_t type.
   QualType OMPAlloctraitT;
   /// Expression for the predefined allocators.
   Expr *OMPPredefinedAllocators[OMPAllocateDeclAttr::OMPUserDefinedMemAlloc] = {
       nullptr};
   /// Vector of previously encountered target directives
   SmallVector<SourceLocation, 2> TargetLocations;
   SourceLocation AtomicLocation;
+  /// Vector of declare variant construct traits.
+  SmallVector<llvm::omp::TraitProperty, 8> ConstructTraits;
 
 public:
   explicit DSAStackTy(Sema &S) : SemaRef(S) {}
 
   /// Sets omp_allocator_handle_t type.
   void setOMPAllocatorHandleT(QualType Ty) { OMPAllocatorHandleT = Ty; }
   /// Gets omp_allocator_handle_t type.
   QualType getOMPAllocatorHandleT() const { return OMPAllocatorHandleT; }
if (VariableCategory == OMPC_DEFAULTMAP_unknown)
                      .ImplicitBehavior != OMPC_DEFAULTMAP_MODIFIER_unknown ||
              getTopOfStack()
                      .DefaultmapMap[OMPC_DEFAULTMAP_pointer]
                      .ImplicitBehavior != OMPC_DEFAULTMAP_MODIFIER_unknown;
     return getTopOfStack().DefaultmapMap[VariableCategory].ImplicitBehavior !=
            OMPC_DEFAULTMAP_MODIFIER_unknown;
   }
 
+  ArrayRef<llvm::omp::TraitProperty> getConstructTraits() {
+    return ConstructTraits;
+  }
+  void handleConstructTrait(ArrayRef<llvm::omp::TraitProperty> Traits,
+                            bool ScopeEntry) {
+    if (ScopeEntry)
+      ConstructTraits.append(Traits.begin(), Traits.end());
+    else
+      for (llvm::omp::TraitProperty Trait : llvm::reverse(Traits)) {
+        llvm::omp::TraitProperty Top = ConstructTraits.pop_back_val();
+        assert(Top == Trait && "Something left a trait on the stack!");
+      }
+  }
+
   DefaultDataSharingAttributes getDefaultDSA(unsigned Level) const {
     return getStackSize() <= Level ? DSA_unspecified
                                    : getStackElemAtLevel(Level).DefaultAttr;
   }
   DefaultDataSharingAttributes getDefaultDSA() const {
     return isStackEmpty() ? DSA_unspecified
                           : getTopOfStack().DefaultAttr;
   }
   SourceLocation getDefaultDSALocation() const {
     return isStackEmpty() ? SourceLocation()
DSAAttrChecker(DSAStackTy *S, Sema &SemaRef, CapturedStmt *CS)
     if (isOpenMPTargetExecutionDirective(S->getCurrentDirective())) {
       for (DeclRefExpr *E : Stack->getLinkGlobals())
         Visit(E);
     }
   }
 };
 } // namespace
 
+static void handleDeclareVariantConstructTrait(DSAStackTy *Stack,
+                                               OpenMPDirectiveKind DKind,
+                                               bool ScopeEntry) {
+  SmallVector<llvm::omp::TraitProperty, 8> Traits;
+  if (isOpenMPTargetExecutionDirective(DKind))
+    Traits.emplace_back(llvm::omp::TraitProperty::construct_target_target);
+  if (isOpenMPTeamsDirective(DKind))
+    Traits.emplace_back(llvm::omp::TraitProperty::construct_teams_teams);
+  if (isOpenMPParallelDirective(DKind))
+    Traits.emplace_back(llvm::omp::TraitProperty::construct_parallel_parallel);
+  if (isOpenMPWorksharingDirective(DKind))
+    Traits.emplace_back(llvm::omp::TraitProperty::construct_for_for);
+  if (isOpenMPSimdDirective(DKind))
+    Traits.emplace_back(llvm::omp::TraitProperty::construct_simd_simd);
+  Stack->handleConstructTrait(Traits, ScopeEntry);
+}
+
 void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   switch (DKind) {
   case OMPD_parallel:
   case OMPD_parallel_for:
   case OMPD_parallel_for_simd:
   case OMPD_parallel_sections:
   case OMPD_parallel_master:
   case OMPD_teams:
void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   case OMPD_begin_declare_variant:
   case OMPD_end_declare_variant:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
   default:
     llvm_unreachable("Unknown OpenMP directive");
   }
   DSAStack->setContext(CurContext);
+  handleDeclareVariantConstructTrait(DSAStack, DKind, /* ScopeEntry */ true);
 }
 
 int Sema::getNumberOfConstructScopes(unsigned Level) const {
   return getOpenMPCaptureLevels(DSAStack->getDirective(Level));
 }
 
 int Sema::getOpenMPCaptureLevels(OpenMPDirectiveKind DKind) {
   SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
S.Diag(Ordered->getBeginLoc(), diag::note_omp_ordered_param)
         << 0 << SourceRange(Ordered->getBeginLoc(), Ordered->getEndLoc());
     return true;
   }
   return false;
 }
 
 StmtResult Sema::ActOnOpenMPRegionEnd(StmtResult S,
                                       ArrayRef<OMPClause *> Clauses) {
+  handleDeclareVariantConstructTrait(DSAStack, DSAStack->getCurrentDirective(),
+                                     /* ScopeEntry */ false);
   if (DSAStack->getCurrentDirective() == OMPD_atomic ||
       DSAStack->getCurrentDirective() == OMPD_critical ||
       DSAStack->getCurrentDirective() == OMPD_section ||
       DSAStack->getCurrentDirective() == OMPD_master ||
       DSAStack->getCurrentDirective() == OMPD_masked)
     return S;
 
   bool ErrorFound = false;
ExprResult Sema::ActOnOpenMPCall(ExprResult Call, Scope *Scope,
   std::function<void(StringRef)> DiagUnknownTrait = [this,
                                                      CE](StringRef ISATrait) {
     // TODO Track the selector locations in a way that is accessible here to
     // improve the diagnostic location.
     Diag(CE->getBeginLoc(), diag::warn_unknown_declare_variant_isa_trait)
         << ISATrait;
   };
   TargetOMPContext OMPCtx(Context, std::move(DiagUnknownTrait),
-                          getCurFunctionDecl());
+                          getCurFunctionDecl(), DSAStack->getConstructTraits());
 
   QualType CalleeFnType = CalleeFnDecl->getType();
 
   SmallVector<Expr *, 4> Exprs;
   SmallVector<VariantMatchInfo, 4> VMIs;
   while (CalleeFnDecl) {
     for (OMPDeclareVariantAttr *A :
          CalleeFnDecl->specific_attrs<OMPDeclareVariantAttr>()) {
       Expr *VariantRef = A->getVariantFuncRef();
 

clang/test/OpenMP/declare_variant_ast_print.c

 // RUN: %clang_cc1 -verify -fopenmp -x c -std=c99 -ast-print %s -o - -Wno-openmp-clauses | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -x c -std=c99 -ast-print %s -o - -Wno-openmp-clauses | FileCheck %s
 
 // expected-no-diagnostics
 
 int foo(void);
 
+#pragma omp declare variant(foo) match(construct={target})
+#pragma omp declare variant(foo) match(construct={teams})
+#pragma omp declare variant(foo) match(construct={parallel})
+#pragma omp declare variant(foo) match(construct={for})
+#pragma omp declare variant(foo) match(construct={simd})
+#pragma omp declare variant(foo) match(construct={target,teams,parallel,for,simd})
 #pragma omp declare variant(foo) match(xxx={}, yyy={ccc})
 #pragma omp declare variant(foo) match(xxx={vvv})
 #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)}, device={kind(fpga)})
 #pragma omp declare variant(foo) match(implementation={vendor(llvm), xxx})
 #pragma omp declare variant(foo) match(implementation={vendor(unknown)}, device={kind(gpu)})
 #pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)}, device={kind(cpu, nohost)})
 #pragma omp declare variant(foo) match(device={kind(host)})
 #pragma omp declare variant(foo) match(device={kind(nohost), xxx})
 #pragma omp declare variant(foo) match(implementation={extension(match_all)})
 #pragma omp declare variant(foo) match(implementation={extension(match_any)})
 #pragma omp declare variant(foo) match(implementation={extension(match_none)})
 int bar(void);
 
 // CHECK:      int foo();
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={extension(match_none)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={extension(match_any)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={extension(match_all)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(nohost)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(host)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm)}, device={kind(cpu, nohost)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(unknown)}, device={kind(gpu)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(llvm)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0): llvm)}, device={kind(fpga)})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(construct={target, teams, parallel, for, simd})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(construct={simd})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(construct={for})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(construct={parallel})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(construct={teams})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(construct={target})
 // CHECK-NEXT: int bar();

clang/test/OpenMP/declare_variant_ast_print.cpp

 // CHECK-NEXT: }
 template <typename T>
 T foofoo() { return T(); }
 
 // CHECK:      template<> int foofoo<int>() {
 // CHECK-NEXT: return int();
 // CHECK-NEXT: }
 
-// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga)})
+// CHECK:      #pragma omp declare variant(foofoo<int>) match(construct={target})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(construct={simd})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga)})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(unknown)})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0): llvm)}, device={kind(cpu)})
 // CHECK-NEXT: int bar();
 #pragma omp declare variant(foofoo <int>) match(xxx = {})
 #pragma omp declare variant(foofoo <int>) match(xxx = {vvv})
 #pragma omp declare variant(foofoo <int>) match(implementation = {vendor(score(0): "llvm"), xxx}, device = {kind(cpu)})
 #pragma omp declare variant(foofoo <int>) match(implementation = {vendor("unknown")})
 #pragma omp declare variant(foofoo <int>) match(implementation = {vendor(score(5): ibm)}, device = {kind(fpga)})
+#pragma omp declare variant(foofoo <int>) match(construct = {simd})
+#pragma omp declare variant(foofoo <int>) match(construct = {target})
 int bar();
 
 // CHECK:      #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5): ibm)}, device={kind(cpu, host)})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(unknown)})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(llvm)}, device={kind(cpu)})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(user={condition(false)})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(user={condition(true)})
 // CHECK-NEXT: template <typename T, int C> T barbar();

clang/test/OpenMP/declare_variant_construct_codegen_1.c

+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK1
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+
+#ifdef CK1
+
+#define N 100
+
+void p_vxv(int *v1, int *v2, int *v3, int n);
+void t_vxv(int *v1, int *v2, int *v3, int n);
+
+#pragma omp declare variant(t_vxv) match(construct={target})
+#pragma omp declare variant(p_vxv) match(construct={parallel})
+void vxv(int *v1, int *v2, int *v3, int n) {
+    for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i];
+}
+// CK1: define dso_local void @vxv
+
+void p_vxv(int *v1, int *v2, int *v3, int n) {
+#pragma omp for
+    for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 3;
+}
+// CK1: define dso_local void @p_vxv
+
+#pragma omp declare target
+void t_vxv(int *v1, int *v2, int *v3, int n) {
+#pragma distribute simd
+    for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 2;
+}
+#pragma omp end declare target
+// CK1: define dso_local void @t_vxv
+
+
+// CK1-LABEL: define {{[^@]+}}@test
+int test() {
+  int v1[N], v2[N], v3[N];
+
+  // init
+  for (int i = 0; i < N; i++) {
+    v1[i] = (i + 1);
+    v2[i] = -(i + 1);
+    v3[i] = 0;
+  }
+
+#pragma omp target teams map(to: v1[:N],v2[:N]) map(from: v3[:N])
+  {
+    vxv(v1, v2, v3, N);
+  }
+// CK1: call void @__omp_offloading_[[OFFLOAD:.+]]({{.+}})
+
+  vxv(v1, v2, v3, N);
+// CK1: call void @vxv
+
+#pragma omp parallel
+  {
+    vxv(v1, v2, v3, N);
+  }
+// CK1: call void ({{.+}}) @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 3, void ({{.+}})* bitcast (void (i32*, i32*, [100 x i32]*, [100 x i32]*, [100 x i32]*)* [[PARALLEL_REGION:@.+]] to void
+
+  return 0;
+}
+
+// CK1: define internal void @__omp_offloading_[[OFFLOAD]]({{.+}})
+// CK1: call void ({{.+}}) @__kmpc_fork_teams(%struct.ident_t* {{.+}}, i32 3, void ({{.+}})* bitcast (void (i32*, i32*, [100 x i32]*, [100 x i32]*, [100 x i32]*)* [[TARGET_REGION:@.+]] to void
+// CK1: define internal void [[TARGET_REGION]](
+// CK1: call void @t_vxv
+
+// CK1: define internal void [[PARALLEL_REGION]](
+// CK1: call void @p_vxv
+#endif // CK1
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK2
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+
+#ifdef CK2
+
+void test_teams(int ***v1, int ***v2, int ***v3, int n);
+void test_target(int ***v1, int ***v2, int ***v3, int n);
+void test_parallel(int ***v1, int ***v2, int ***v3, int n);
+
+#pragma omp declare variant(test_teams) match(construct = {teams})
+#pragma omp declare variant(test_target) match(construct = {target})
+#pragma omp declare variant(test_parallel) match(construct = {parallel})
+void test_base(int ***v1, int ***v2, int ***v3, int n) {
+  for (int i = 0; i < n; i++)
+    for (int j = 0; j < n; ++j)
+      for (int k = 0; k < n; ++k)
+        v3[i][j][k] = v1[i][j][k] * v2[i][j][k];
+}
+
+#pragma omp declare target
+void test_teams(int ***v1, int ***v2, int ***v3, int n) {
+#pragma omp distribute parallel for simd collapse(2)
+  for (int i = 0; i < n; ++i)
+    for (int j = 0; j < n; ++j)
+      for (int k = 0; k < n; ++k)
+        v3[i][j][k] = v1[i][j][k] * v2[i][j][k];
+}
+#pragma omp end declare target
+
+#pragma omp declare target
+void test_target(int ***v1, int ***v2, int ***v3, int n) {
+#pragma omp parallel for simd collapse(3)
+  for (int i = 0; i < n; ++i)
+    for (int j = 0; j < n; ++j)
+      for (int k = 0; k < n; ++k)
+        v3[i][j][k] = v1[i][j][k] * v2[i][j][k];
+}
+#pragma omp end declare target
+
+void test_parallel(int ***v1, int ***v2, int ***v3, int n) {
+#pragma omp for collapse(3)
+  for (int i = 0; i < n; ++i)
+    for (int j = 0; j < n; ++j)
+      for (int k = 0; k < n; ++k)
+        v3[i][j][k] = v1[i][j][k] * v2[i][j][k];
+}
+
+// CK2-LABEL: define {{[^@]+}}@test
+void test(int ***v1, int ***v2, int ***v3, int n) {
+  int i;
+
+#pragma omp target
+#pragma omp teams
+  {
+    test_base(v1, v2, v3, 0);
+  }
+// CK2: call void @__omp_offloading_[[OFFLOAD_1:.+]]({{.+}})
+
+#pragma omp target
+  {
+    test_base(v1, v2, v3, 0);
+  }
+// CK2: call void @__omp_offloading_[[OFFLOAD_2:.+]]({{.+}})
+
+#pragma omp parallel
+  {
+    test_base(v1, v2, v3, 0);
+  }
+// CK2: call void ({{.+}}) @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 3, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32****, i32****, i32****)* [[PARALLEL_REGION:@.+]] to void
+}
+
+// CK2: define internal void @__omp_offloading_[[OFFLOAD_1]]({{.+}})
+// CK2: call void ({{.+}}) @__kmpc_fork_teams(%struct.ident_t* {{.+}}, i32 3, void ({{.+}})* bitcast (void (i32*, i32*, i32****, i32****, i32****)* [[TARGET_REGION_1:@.+]] to void
+// CK2: define internal void [[TARGET_REGION_1]](
+// CK2: call void @test_teams
+
+// CK2: define internal void @__omp_offloading_[[OFFLOAD_2]]({{.+}})
+// CK2: call void @test_target
+
+// CK2: define internal void [[PARALLEL_REGION]](
+// CK2: call void @test_parallel
+
+#endif // CK2
+
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK3
+// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK3
+// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK3
+// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK3
+// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK3
+
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+
+#ifdef CK3
+
+#define N 100
+
+int t_for(int *v1, int *v2, int *v3, int n);
+int t_simd(int *v1, int *v2, int *v3, int n);
+
+#pragma omp declare variant(t_simd) match(construct = {simd})
+#pragma omp declare variant(t_for) match(construct = {for})
+int t(int *v1, int *v2, int *v3, int idx) {
+  return v1[idx] * v2[idx];
+}
+
+int t_for(int *v1, int *v2, int *v3, int idx) {
+  return v1[idx] * v2[idx];
+}
+
+#pragma omp declare simd
+int t_simd(int *v1, int *v2, int *v3, int idx) {
+  return v1[idx] * v2[idx];
+}
+
+// CK3-LABEL: define {{[^@]+}}@test
+void test() {
+  int v1[N], v2[N], v3[N];
+
+  // init
+  for (int i = 0; i < N; i++) {
+    v1[i] = (i + 1);
+    v2[i] = -(i + 1);
+    v3[i] = 0;
+  }
+
+#pragma omp simd
+  for (int i = 0; i < N; i++) {
+    v3[i] = t(v1, v2, v3, i);
+  }
+// CK3: call = call i32 @t_simd
+
+
+#pragma omp for
+  for (int i = 0; i < N; i++) {
+    v3[i] = t(v1, v2, v3, i);
+  }
+// CK3: call{{.+}} = call i32 @t_for
+}
+
+#endif // CK3
+
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK4
+// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK4
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK4
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK4
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK4
+
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
+
+#ifdef CK4
+
+#define N 100
+
+void not_selected_vxv(int *v1, int *v2, int *v3, int n);
+void combined_vxv(int *v1, int *v2, int *v3, int n);
+void all_vxv(int *v1, int *v2, int *v3, int n);
+
+#pragma omp declare variant(all_vxv) match(construct={target,teams,parallel,for,simd})
+#pragma omp declare variant(combined_vxv) match(construct={target,teams,parallel,for})
+#pragma omp declare variant(not_selected_vxv) match(construct={parallel,for})
+void vxv(int *v1, int *v2, int *v3, int n) {
+    for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i];
+}
+
+void not_selected_vxv(int *v1, int *v2, int *v3, int n) {
+    for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 3;
+}
+
+#pragma omp declare target
+void combined_vxv(int *v1, int *v2, int *v3, int n) {
+    for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 2;
+}
+#pragma omp end declare target
+
+#pragma omp declare target
+void all_vxv(int *v1, int *v2, int *v3, int n) {
+    for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 4;
+}
+#pragma omp end declare target
+
+// CK4-LABEL: define {{[^@]+}}@test
+void test() {
+    int v1[N], v2[N], v3[N];
+
+    //init
+    for (int i = 0; i < N; i++) {
+      v1[i] = (i + 1);
+      v2[i] = -(i + 1);
+      v3[i] = 0;
+    }
+
+#pragma omp target teams map(to: v1[:N],v2[:N]) map(from: v3[:N])
+    {
+#pragma omp parallel for
+      for (int i = 0; i < N; i++)
+        vxv(v1, v2, v3, N);
+    }
+// CK4: call void @__omp_offloading_[[OFFLOAD_1:.+]]({{.+}})
+
+#pragma omp simd
+    for (int i = 0; i < N; i++)
+      vxv(v1, v2, v3, N);
+// CK4: call void @vxv
+
+#pragma omp target teams distribute parallel for simd map(from: v3[:N])
+    for (int i = 0; i < N; i++)
+      for (int i = 0; i < N; i++)
+        for (int i = 0; i < N; i++)
+          vxv(v1, v2, v3, N);
+// CK4: call void @__omp_offloading_[[OFFLOAD_2:.+]]({{.+}})
+}
+// CK4-DAG: call void @all_vxv
+// CK4-DAG: call void @combined_vxv
+
+#endif // CK4
+
+#endif // HEADER

clang/test/OpenMP/declare_variant_messages.c

 #pragma omp declare variant(foo) match(implementation = {vendor(score(foo) :llvm)}) // expected-warning {{score expressions in the OpenMP context selector need to be constant; foo is not and will be ignored}}
 #pragma omp declare variant(foo) match(implementation = {vendor(score(foo()) :llvm)}) // expected-warning {{score expressions in the OpenMP context selector need to be constant; foo() is not and will be ignored}}
 #pragma omp declare variant(foo) match(implementation = {vendor(score(<expr>) :llvm)}) // expected-error {{expected expression}} expected-error {{use of undeclared identifier 'expr'}} expected-error {{expected expression}}
 #pragma omp declare variant(foo) match(user = {condition(foo)}) // expected-error {{the user condition in the OpenMP context selector needs to be constant; foo is not}}
 #pragma omp declare variant(foo) match(user = {condition(foo())}) // expected-error {{the user condition in the OpenMP context selector needs to be constant; foo() is not}}
 #pragma omp declare variant(foo) match(user = {condition(<expr>)}) // expected-error {{expected expression}} expected-error {{use of undeclared identifier 'expr'}} expected-error {{expected expression}} expected-note {{the ignored selector spans until here}}
 int score_and_cond_non_const();
 
+#pragma omp declare variant(foo) match(construct={teams,parallel,for,simd})
+#pragma omp declare variant(foo) match(construct={target teams}) // expected-error {{expected ')'}} expected-warning {{expected '}' after the context selectors for the context set "construct"; '}' assumed}} expected-note {{to match this '('}}
+#pragma omp declare variant(foo) match(construct={parallel for}) // expected-error {{expected ')'}} expected-warning {{expected '}' after the context selectors for the context set "construct"; '}' assumed}} expected-note {{to match this '('}}
+#pragma omp declare variant(foo) match(construct={for simd}) // expected-error {{expected ')'}} expected-warning {{expected '}' after the context selectors for the context set "construct"; '}' assumed}} expected-note {{to match this '('}}
+int construct(void);
+
 #pragma omp declare variant(foo) match(xxx={}) // expected-warning {{'xxx' is not a valid context set in a `declare variant`; set ignored}} expected-note {{context set options are: 'construct' 'device' 'implementation' 'user'}} expected-note {{the ignored set spans until here}}
 int a; // expected-error {{'#pragma omp declare variant' can only be applied to functions}}
 
 #pragma omp declare variant(foo) match(xxx={}) // expected-warning {{'xxx' is not a valid context set in a `declare variant`; set ignored}} expected-note {{context set options are: 'construct' 'device' 'implementation' 'user'}} expected-note {{the ignored set spans until here}}
 #pragma omp threadprivate(a) // expected-error {{'#pragma omp declare variant' can only be applied to functions}}
 int var;
 #pragma omp threadprivate(var)