https://github.com/chichunchen updated https://github.com/llvm/llvm-project/pull/231668
>From 700ebcef8d8f866b31299085a223a32bc9eaf05c Mon Sep 17 00:00:00 2001 From: "Chi Chun, Chen" <[email protected]> Date: Fri, 16 Oct 2026 12:22:52 -0500 Subject: [PATCH 1/2] [OpenMP] Fix context selector matching and scoring OMPContext computes some context selector scores differently from the OpenMP 5.2 rules in "Matching and Scoring Context Selectors", and crashes on some selectors. Fix the following; each can change which variant is selected: - Device weights. kind, arch, and isa are worth 2^l, 2^(l+1), and 2^(l+2), where l is the depth of the construct context. OMPContext used the candidate's own construct selector count, so inside two nested parallel regions kind(cpu) was worth 1 instead of 4 and lost to vendor(score(2): llvm). - Scores per selector. A score belongs to a trait selector, but it was added once per property, so vendor(score(10): llvm, amd) counted 20. Scores now add once, using arbitrary-width arithmetic instead of wrapping 64-bit sums. - Repeated constructs. When a construct selector can match several enclosing constructs, the highest-valued complete match counts. construct={parallel} inside three nested parallel regions now adds 2^2 instead of 2^0. - match_any and match_none with construct selectors. match_any stopped at the first matching property before recording construct positions, which tripped an assertion when the variant was scored. match_none rejected a construct selector even outside that construct. - Unknown properties. Flang aborted with "Unknown trait set is not to be used!" when scoring an unknown property such as vendor(bogus). Such a property now adds nothing unless it has a user score. Left for follow-ups: - The strict-subset rule still only breaks ties. Applying it to every candidate requires comparing user conditions. - Flang records an unknown property as invalid, without its selector. An unknown kind or arch property therefore gets no 2^l weight, and a user score on a selector that mixes known and unknown properties counts twice. Clang drops unknown properties while parsing. - Under match_any and match_none, a selector adds its weight even when it does not match, as before. With the fixes above, an unmatched kind(gpu) can now tie a matching kind(host, cpu) or outweigh a vendor score. - Clang's construct context is incomplete. Declare variant counts single and sections as for, which now also raises device weights there. Metadirectives see no enclosing constructs, so match_none with a construct selector now matches even inside that construct. Split from #224431. Assisted-by: GitHub Copilot (Claude Opus 5.5) --- .../declare_variant_construct_match_kind.c | 30 ++ .../OpenMP/metadirective-implementation.f90 | 12 + llvm/lib/Frontend/OpenMP/OMPContext.cpp | 196 ++++++------- llvm/unittests/Frontend/OpenMPContextTest.cpp | 258 +++++++++++++++++- 4 files changed, 391 insertions(+), 105 deletions(-) create mode 100644 clang/test/OpenMP/declare_variant_construct_match_kind.c diff --git a/clang/test/OpenMP/declare_variant_construct_match_kind.c b/clang/test/OpenMP/declare_variant_construct_match_kind.c new file mode 100644 index 00000000000000..db6a475a5bd7a3 --- /dev/null +++ b/clang/test/OpenMP/declare_variant_construct_match_kind.c @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - \ +// RUN: | FileCheck %s + +int v1(void); +int v2(void); + +// MATCH_NONE accepts a construct selector outside that construct. +#pragma omp declare variant(v1) \ + match(construct = {parallel}, implementation = {extension(match_none)}) +int none_base(void); + +// CHECK-LABEL: define{{.*}} i32 @none_caller( +// CHECK: call i32 @v1() +int none_caller(void) { return none_base(); } + +// Under MATCH_ANY, the matching construct adds its weight: 2 against 1. +#pragma omp declare variant(v1) \ + match(construct = {parallel}, \ + implementation = {vendor(llvm), extension(match_any)}) +#pragma omp declare variant(v2) match(user = {condition(1)}) +int any_base(void); + +// CHECK-LABEL: define{{.*}} @any_caller.omp_outlined( +// CHECK: call i32 @v1() +int any_caller(void) { + int r = 0; +#pragma omp parallel + r = any_base(); + return r; +} diff --git a/flang/test/Lower/OpenMP/metadirective-implementation.f90 b/flang/test/Lower/OpenMP/metadirective-implementation.f90 index af8f2af938ca55..cd61b733cbb1d3 100644 --- a/flang/test/Lower/OpenMP/metadirective-implementation.f90 +++ b/flang/test/Lower/OpenMP/metadirective-implementation.f90 @@ -5,6 +5,18 @@ ! RUN: %flang_fc1 -fopenmp -emit-hlfir -fopenmp-version=51 %s -o - | FileCheck %s ! RUN: %flang_fc1 -fopenmp -emit-hlfir -fopenmp-version=52 -cpp -DOMP_52 %s -o - | FileCheck %s +! MATCH_NONE accepts an unknown vendor; clause order breaks the scoring tie. +! CHECK-LABEL: func.func @_QPtest_unknown_vendor_match_none() +! CHECK-NOT: omp.barrier +! CHECK: omp.taskyield +! CHECK-NEXT: return +subroutine test_unknown_vendor_match_none() + !$omp metadirective & + !$omp& when(implementation={vendor(bogus_vendor), extension(match_none)}: & + !$omp& taskyield) & + !$omp& when(user={condition(.true.)}: barrier) +end subroutine + ! CHECK-LABEL: func.func @_QPtest_vendor_llvm() ! CHECK: omp.taskwait ! CHECK: return diff --git a/llvm/lib/Frontend/OpenMP/OMPContext.cpp b/llvm/lib/Frontend/OpenMP/OMPContext.cpp index e77e48913fdd0c..b87bc7a93bd4f0 100644 --- a/llvm/lib/Frontend/OpenMP/OMPContext.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPContext.cpp @@ -13,11 +13,13 @@ //===----------------------------------------------------------------------===// #include "llvm/Frontend/OpenMP/OMPContext.h" +#include "llvm/ADT/DenseSet.h" #include "llvm/ADT/StringRef.h" #include "llvm/ADT/StringSwitch.h" #include "llvm/Support/Debug.h" #include "llvm/Support/raw_ostream.h" #include "llvm/TargetParser/Triple.h" +#include <algorithm> #define DEBUG_TYPE "openmp-ir-builder" @@ -207,23 +209,21 @@ isVariantApplicableInContextHelper(const VariantMatchInfo &VMI, unsigned(TraitProperty::implementation_extension_match_none))) MK = MK_NONE; - // Helper to deal with a single property that was (not) found in the OpenMP - // context based on the match kind selected by the user via - // `implementation={extensions(match_[all,any,none])}' - auto HandleTrait = [MK](TraitProperty Property, - bool WasFound) -> std::optional<bool> /* Result */ { - // For kind "any" a single match is enough but we ignore non-matched - // properties. - if (MK == MK_ANY) { - if (WasFound) - return true; - return std::nullopt; - } + bool AnyTraitMatched = false; + + // Apply the match kind selected by implementation={extension(...)} to + // each property. Continue after match_any succeeds to record all construct + // match positions needed for scoring. + auto HandleTrait = [MK, &AnyTraitMatched](TraitProperty Property, + bool WasFound) -> bool { + AnyTraitMatched |= WasFound; + if (MK == MK_ANY) + return true; // In "all" or "none" mode we accept a matching or non-matching property // respectively and move on. We are not done yet! if ((WasFound && MK == MK_ALL) || (!WasFound && MK == MK_NONE)) - return std::nullopt; + return true; // We missed a property, provide some debug output and indicate failure. LLVM_DEBUG({ @@ -266,30 +266,35 @@ isVariantApplicableInContextHelper(const VariantMatchInfo &VMI, return Ctx.matchesISATrait(RawString); }); - if (std::optional<bool> Result = HandleTrait(Property, IsActiveTrait)) - return *Result; + if (!HandleTrait(Property, IsActiveTrait)) + return false; + // The match_any result is known; skip ISA lookups, which can warn. + if (MK == MK_ANY && AnyTraitMatched) + break; } if (!DeviceOrImplementationSetOnly) { - // We could use isOrderedSubset here but we also want to record the match - // locations. + // Scan the construct sequence in order, recording matching context + // positions for scoring. unsigned ConstructIdx = 0, NoConstructTraits = Ctx.ConstructTraits.size(); for (TraitProperty Property : VMI.ConstructTraits) { assert(getOpenMPContextTraitSetForProperty(Property) == TraitSet::construct && "Variant context is ill-formed!"); - // Verify the nesting. + // Verify the nesting. A failed match in match_any or match_none must not + // consume the remaining context, since a later selector property can + // still match. + unsigned SearchStart = ConstructIdx; bool FoundInOrder = false; while (!FoundInOrder && ConstructIdx != NoConstructTraits) FoundInOrder = (Ctx.ConstructTraits[ConstructIdx++] == Property); - if (ConstructMatches) + if (!FoundInOrder && MK != MK_ALL) + ConstructIdx = SearchStart; + if (ConstructMatches && FoundInOrder) ConstructMatches->push_back(ConstructIdx - 1); - if (std::optional<bool> Result = HandleTrait(Property, FoundInOrder)) - return *Result; - - if (!FoundInOrder) { + if (!HandleTrait(Property, FoundInOrder)) { LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] Construct property " << getOpenMPContextTraitPropertyName(Property, "") << " was not nested properly.\n"); @@ -299,12 +304,29 @@ isVariantApplicableInContextHelper(const VariantMatchInfo &VMI, // TODO: Verify SIMD } - assert(isOrderedSubset<TraitProperty>(VMI.ConstructTraits, - Ctx.ConstructTraits) && - "Broken invariant!"); + // A complete ordered match can have several embeddings in the context. + // Match backwards to choose the highest-valued one for scoring. Keep the + // forward scan's partial matches for the match_any extension. + if (ConstructMatches && + ConstructMatches->size() == VMI.ConstructTraits.size()) { + ConstructIdx = NoConstructTraits; + for (unsigned I = VMI.ConstructTraits.size(); I > 0; --I) { + TraitProperty Property = VMI.ConstructTraits[I - 1]; + while (ConstructIdx > 0 && + Ctx.ConstructTraits[ConstructIdx - 1] != Property) + --ConstructIdx; + assert(ConstructIdx > 0 && "Previously matched construct not found!"); + (*ConstructMatches)[I - 1] = --ConstructIdx; + } + } + + if (MK == MK_ALL) + assert(isOrderedSubset<TraitProperty>(VMI.ConstructTraits, + Ctx.ConstructTraits) && + "Broken invariant!"); } - if (MK == MK_ANY) { + if (MK == MK_ANY && !AnyTraitMatched) { LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] None of the properties was in the OpenMP context " "but match kind is any.\n"); @@ -324,91 +346,78 @@ bool llvm::omp::isVariantApplicableInContext( static APInt getVariantMatchScore(const VariantMatchInfo &VMI, const OMPContext &Ctx, SmallVectorImpl<unsigned> &ConstructMatches) { - APInt Score(64, 1); - - unsigned NoConstructTraits = VMI.ConstructTraits.size(); - for (unsigned Bit : VMI.RequiredTraits.set_bits()) { - TraitProperty Property = TraitProperty(Bit); - // If there is a user score attached, use it. - if (VMI.ScoreMap.count(Property)) { - const APInt &UserScore = VMI.ScoreMap.lookup(Property); - assert(UserScore.uge(0) && "Expect non-negative user scores!"); - Score += UserScore.getZExtValue(); - continue; - } + APInt Score(1, 1); + + // A sum of valid scores can exceed the width of any individual score. + // Retain all active bits and allow one more bit for each addition's carry. + auto AddScore = [&](const APInt &Value) { + unsigned Width = std::max(Score.getActiveBits(), Value.getActiveBits()) + 1; + Score = Score.zextOrTrunc(Width); + Score += Value.zextOrTrunc(Width); + }; + auto AddPowerOfTwo = [&](unsigned Exponent) { + AddScore(APInt::getOneBitSet(Exponent + 1, Exponent)); + }; - switch (getOpenMPContextTraitSetForProperty(Property)) { - case TraitSet::construct: - // We handle the construct traits later via the VMI.ConstructTraits - // container. - continue; - case TraitSet::implementation: - // No effect on the score (implementation defined). - continue; - case TraitSet::user: - // No effect on the score. - continue; - case TraitSet::device: - // Handled separately below. - break; - case TraitSet::target_device: - // TODO: Handling separately. - break; - case TraitSet::invalid: - llvm_unreachable("Unknown trait set is not to be used!"); + unsigned NoConstructTraits = Ctx.ConstructTraits.size(); + SmallDenseSet<TraitSelector, 8> ScoredSelectors; + auto AddSelectorScore = [&](TraitSelector Selector, const APInt *UserScore) { + // Scores belong to selectors, not to individual properties. + if (!ScoredSelectors.insert(Selector).second) + return; + if (UserScore) { + AddScore(*UserScore); + return; } - - // device={kind(any)} is "as if" no kind selector was specified. - if (Property == TraitProperty::device_kind_any) - continue; - if (Property == TraitProperty::target_device_kind_any) - continue; - - switch (getOpenMPContextTraitSelectorForProperty(Property)) { + switch (Selector) { case TraitSelector::device_kind: - Score += (1ULL << (NoConstructTraits + 0)); - continue; - case TraitSelector::device_arch: - Score += (1ULL << (NoConstructTraits + 1)); - continue; - case TraitSelector::device_isa: - Score += (1ULL << (NoConstructTraits + 2)); - continue; case TraitSelector::target_device_kind: - Score += (1ULL << (NoConstructTraits + 0)); - continue; + AddPowerOfTwo(NoConstructTraits); + break; + case TraitSelector::device_arch: case TraitSelector::target_device_arch: - Score += (1ULL << (NoConstructTraits + 1)); - continue; + AddPowerOfTwo(NoConstructTraits + 1); + break; + case TraitSelector::device_isa: case TraitSelector::target_device_isa: - Score += (1ULL << (NoConstructTraits + 2)); - continue; + AddPowerOfTwo(NoConstructTraits + 2); + break; default: - continue; + break; } + }; + for (unsigned Bit : VMI.RequiredTraits.set_bits()) { + TraitProperty Property = TraitProperty(Bit); + // Construct scores use ordered positions below. kind(any) is treated as + // if no kind selector were specified. + if (getOpenMPContextTraitSetForProperty(Property) == TraitSet::construct || + Property == TraitProperty::device_kind_any || + Property == TraitProperty::target_device_kind_any) + continue; + auto It = VMI.ScoreMap.find(Property); + AddSelectorScore(getOpenMPContextTraitSelectorForProperty(Property), + It == VMI.ScoreMap.end() ? nullptr : &It->second); } - unsigned ConstructIdx = 0; - assert(NoConstructTraits == ConstructMatches.size() && + assert(VMI.ConstructTraits.size() >= ConstructMatches.size() && "Mismatch in the construct traits!"); - for (TraitProperty Property : VMI.ConstructTraits) { - assert(getOpenMPContextTraitSetForProperty(Property) == - TraitSet::construct && - "Ill-formed variant match info!"); - (void)Property; + for (unsigned Match : ConstructMatches) { // ConstructMatches is the position p - 1 and we need 2^(p-1). - Score += (1ULL << ConstructMatches[ConstructIdx++]); + AddPowerOfTwo(Match); } - LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] Variant has a score of " << Score - << "\n"); + LLVM_DEBUG({ + dbgs() << "[" << DEBUG_TYPE << "] Variant has a score of "; + Score.print(dbgs(), /*isSigned=*/false); + dbgs() << "\n"; + }); return Score; } int llvm::omp::getBestVariantMatchForContext( const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx) { - APInt BestScore(64, 0); + APInt BestScore(1, 0); int BestVMIIdx = -1; const VariantMatchInfo *BestVMI = nullptr; @@ -423,6 +432,9 @@ int llvm::omp::getBestVariantMatchForContext( continue; // Check if its clearly not the best. APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches); + unsigned Width = std::max(Score.getBitWidth(), BestScore.getBitWidth()); + Score = Score.zextOrTrunc(Width); + BestScore = BestScore.zextOrTrunc(Width); if (Score.ult(BestScore)) continue; // Equal score need subset checks. diff --git a/llvm/unittests/Frontend/OpenMPContextTest.cpp b/llvm/unittests/Frontend/OpenMPContextTest.cpp index ddba513e5330cb..2179193df5948d 100644 --- a/llvm/unittests/Frontend/OpenMPContextTest.cpp +++ b/llvm/unittests/Frontend/OpenMPContextTest.cpp @@ -314,7 +314,242 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) { } TEST_F(OpenMPContextTest, ScoringSimple) { - // TODO: Add scoring tests (via getBestVariantMatchForContext). + OMPContext Parallel(false, Triple("x86_64-unknown-linux"), Triple(), -1); + Parallel.addTrait(TraitProperty::construct_parallel_parallel); + OMPContext NoConstruct(false, Triple("x86_64-unknown-linux"), Triple(), -1); + + VariantMatchInfo MatchAny; + MatchAny.addTrait(TraitProperty::construct_target_target, ""); + MatchAny.addTrait(TraitProperty::construct_parallel_parallel, ""); + MatchAny.addTrait(TraitProperty::implementation_extension_match_any, ""); + EXPECT_TRUE(isVariantApplicableInContext(MatchAny, Parallel)); + EXPECT_FALSE(isVariantApplicableInContext(MatchAny, NoConstruct)); + + VariantMatchInfo VendorLLVM; + VendorLLVM.addTrait(TraitProperty::implementation_vendor_llvm, ""); + // The matching construct must raise the score, not just win a tie by order. + SmallVector<VariantMatchInfo, 2> MatchAnyCandidates{VendorLLVM, MatchAny}; + EXPECT_EQ(getBestVariantMatchForContext(MatchAnyCandidates, Parallel), 1); + + VariantMatchInfo MatchNone; + MatchNone.addTrait(TraitProperty::construct_parallel_parallel, ""); + MatchNone.addTrait(TraitProperty::implementation_extension_match_none, ""); + EXPECT_TRUE(isVariantApplicableInContext(MatchNone, NoConstruct)); + EXPECT_FALSE(isVariantApplicableInContext(MatchNone, Parallel)); + + VariantMatchInfo Empty; + SmallVector<VariantMatchInfo, 2> MatchNoneCandidates{MatchNone, Empty}; + EXPECT_EQ(getBestVariantMatchForContext(MatchNoneCandidates, NoConstruct), 0); +} + +TEST_F(OpenMPContextTest, ScoringMatchAnyConstructs) { + OMPContext TargetParallel(false, Triple("x86_64-unknown-linux"), Triple(), + -1); + TargetParallel.addTrait(TraitProperty::construct_target_target); + TargetParallel.addTrait(TraitProperty::construct_parallel_parallel); + + APInt Score(64, 2); + VariantMatchInfo Scored; + Scored.addTrait(TraitProperty::user_condition_true, "", &Score); + + VariantMatchInfo MatchAny; + MatchAny.addTrait(TraitProperty::construct_target_target, ""); + MatchAny.addTrait(TraitProperty::construct_parallel_parallel, ""); + MatchAny.addTrait(TraitProperty::implementation_extension_match_any, ""); + + // Both construct matches contribute: 1 + 1 + 2 for MatchAny beats Scored's + // 1 + 2. Keeping only one of the construct matches would not select MatchAny. + SmallVector<VariantMatchInfo, 2> Candidates{Scored, MatchAny}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, TargetParallel), 1); +} + +TEST_F(OpenMPContextTest, ScoringMatchAnyWithoutMatchingConstructs) { + OMPContext NoConstruct(false, Triple("x86_64-unknown-linux"), Triple(), -1); + + VariantMatchInfo MatchAny; + MatchAny.addTrait(TraitProperty::construct_parallel_parallel, ""); + MatchAny.addTrait(TraitProperty::implementation_vendor_llvm, ""); + MatchAny.addTrait(TraitProperty::implementation_extension_match_any, ""); + EXPECT_TRUE(isVariantApplicableInContext(MatchAny, NoConstruct)); + + APInt Score(64, 1); + VariantMatchInfo Scored; + Scored.addTrait(TraitProperty::user_condition_true, "", &Score); + + // The vendor match makes MatchAny applicable, but the absent construct adds + // no score. Scored wins by 2 to 1. + SmallVector<VariantMatchInfo, 2> Candidates{MatchAny, Scored}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, NoConstruct), 1); +} + +TEST_F(OpenMPContextTest, ScoringInvalidProperty) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + VariantMatchInfo Unknown; + Unknown.addTrait(TraitProperty::invalid, "bogus_vendor"); + EXPECT_FALSE(isVariantApplicableInContext(Unknown, Context)); + Unknown.addTrait(TraitProperty::implementation_extension_match_none, ""); + EXPECT_TRUE(isVariantApplicableInContext(Unknown, Context)); + + VariantMatchInfo UserTrue; + UserTrue.addTrait(TraitProperty::user_condition_true, ""); + SmallVector<VariantMatchInfo, 2> Candidates{Unknown, UserTrue}; + // Scoring an unknown property must not abort. Both score 1, so candidate + // order breaks the tie. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + + APInt Score(64, 1); + Candidates[1].addTrait(TraitProperty::user_condition_true, "", &Score); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); +} + +TEST_F(OpenMPContextTest, ScoringUserScoreOncePerSelector) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + APInt VendorScore(64, 10), Score(64, 11); + VariantMatchInfo TwoVendors, Scored; + TwoVendors.addTrait(TraitProperty::implementation_vendor_llvm, "", + &VendorScore); + TwoVendors.addTrait(TraitProperty::implementation_vendor_amd, "", + &VendorScore); + TwoVendors.addTrait(TraitProperty::implementation_extension_match_any, ""); + Scored.addTrait(TraitProperty::user_condition_true, "", &Score); + + // The vendor score counts once: 1 + 10 loses to 1 + 11. + SmallVector<VariantMatchInfo, 2> Candidates{TwoVendors, Scored}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); +} + +TEST_F(OpenMPContextTest, ScoringMultipleProperties) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + VariantMatchInfo Device; + Device.addTrait(TraitProperty::device_kind_cpu, ""); + Device.addTrait(TraitProperty::device_kind_host, ""); + APInt Score(64, 1); + VariantMatchInfo Scored; + Scored.addTrait(TraitProperty::implementation_vendor_llvm, "", &Score); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, Device}; + // KIND contributes one weight regardless of its number of properties. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); +} + +TEST_F(OpenMPContextTest, ScoringRepeatedConstructs) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + Context.addTrait(TraitProperty::construct_parallel_parallel); + Context.addTrait(TraitProperty::construct_parallel_parallel); + + APInt Score(64, 1); + VariantMatchInfo Scored; + Scored.addTrait(TraitProperty::user_condition_true, "", &Score); + VariantMatchInfo Parallel; + Parallel.addTrait(TraitProperty::construct_parallel_parallel, ""); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, Parallel}; + // The inner PARALLEL scores 3, beating the explicit score's 2. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + Candidates[1].addTrait(TraitProperty::implementation_extension_match_any, ""); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + + // An incomplete match_any selector retains its forward partial match. + VariantMatchInfo Partial; + Partial.addTrait(TraitProperty::construct_target_target, ""); + Partial.addTrait(TraitProperty::construct_parallel_parallel, ""); + Partial.addTrait(TraitProperty::implementation_extension_match_any, ""); + Candidates[1] = Partial; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); +} + +TEST_F(OpenMPContextTest, ScoringHighestOrderedMatch) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + Context.addTrait(TraitProperty::construct_parallel_parallel); + Context.addTrait(TraitProperty::construct_for_for); + Context.addTrait(TraitProperty::construct_parallel_parallel); + Context.addTrait(TraitProperty::construct_for_for); + Context.addTrait(TraitProperty::construct_parallel_parallel); + + APInt Score(64, 11); + VariantMatchInfo Scored; + Scored.addTrait(TraitProperty::user_condition_true, "", &Score); + VariantMatchInfo ParallelFor; + ParallelFor.addTrait(TraitProperty::construct_parallel_parallel, ""); + ParallelFor.addTrait(TraitProperty::construct_for_for, ""); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, ParallelFor}; + // Positions 3 and 4 score 1 + 4 + 8 = 13. The final PARALLEL cannot be + // chosen because it follows every FOR in the context. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + Score = APInt(64, 12); + Candidates[0].addTrait(TraitProperty::user_condition_true, "", &Score); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); +} + +TEST_F(OpenMPContextTest, ScoringDeviceWeights) { + struct ISAContext : OMPContext { + using OMPContext::OMPContext; + bool matchesISATrait(StringRef) const override { return true; } + }; + const TraitProperty Properties[] = {TraitProperty::device_kind_cpu, + TraitProperty::device_arch_arm, + TraitProperty::device_isa___ANY, + TraitProperty::target_device_kind_cpu, + TraitProperty::target_device_arch_arm, + TraitProperty::target_device_isa___ANY}; + for (unsigned Depth : {0u, 1u, 2u, 63u, 64u, 65u}) { + ISAContext Context(false, Triple("arm-unknown-linux"), Triple(), -1); + for (unsigned I = 0; I < Depth; ++I) + Context.addTrait(TraitProperty::construct_parallel_parallel); + for (unsigned I = 0; I < 6; ++I) { + SCOPED_TRACE(Depth); + SCOPED_TRACE(I); + APInt Weight = APInt::getOneBitSet(128, Depth + I % 3); + APInt Score = Weight - 1; + VariantMatchInfo Scored, Device; + Scored.addTrait(TraitProperty::user_condition_true, "", &Score); + Device.addTrait(Properties[I], "test-isa"); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, Device}; + // The device score is exactly 1 + 2^(context depth + trait offset). + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + Candidates[0].addTrait(TraitProperty::user_condition_true, "", &Weight); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + } + } +} + +TEST_F(OpenMPContextTest, ScoringWideTotals) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + Context.addTrait(TraitProperty::construct_parallel_parallel); + for (unsigned Bits : {63u, 64u, 127u}) { + SCOPED_TRACE(Bits); + APInt Score = APInt::getLowBitsSet(Bits + 1, Bits); + VariantMatchInfo Large, Small; + Large.addTrait(TraitProperty::user_condition_true, "", &Score); + Large.addTrait(TraitProperty::implementation_vendor_llvm, "", &Score); + Large.addTrait(TraitProperty::construct_parallel_parallel, ""); + APInt SmallScore(32, 20); + Small.addTrait(TraitProperty::user_condition_true, "", &SmallScore); + SmallVector<VariantMatchInfo, 2> Candidates{Large, Small}; + // The sum is 2^(Bits + 1), not zero after fixed-width overflow. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + std::swap(Candidates[0], Candidates[1]); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + // Overflow to a nonzero value must not silently reverse the ranking. + Candidates[1].addTrait(TraitProperty::device_kind_cpu, ""); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + } +} + +TEST_F(OpenMPContextTest, ScoringDeepConstructs) { + for (unsigned Depth : {63u, 64u, 65u}) { + SCOPED_TRACE(Depth); + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + for (unsigned I = 0; I < Depth; ++I) + Context.addTrait(TraitProperty::construct_parallel_parallel); + APInt Weight = APInt::getOneBitSet(128, Depth - 1); + APInt Score = Weight - 1; + VariantMatchInfo Scored, Parallel; + Scored.addTrait(TraitProperty::user_condition_true, "", &Score); + Parallel.addTrait(TraitProperty::construct_parallel_parallel, ""); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, Parallel}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + Candidates[0].addTrait(TraitProperty::user_condition_true, "", &Weight); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + } } TEST_F(OpenMPContextTest, ConstructContextStartsAtTarget) { @@ -391,17 +626,16 @@ TEST_F(OpenMPContextTest, RepeatedConstructWithConditionOrISA) { // Each pair ties and requires the same traits, but the conditions or ISA // names may differ, so clause order decides. - APInt Four(64, 4), Twenty(64, 20); - for (auto [Property, Score] : - {std::pair(TraitProperty::user_condition_true, &Four), - std::pair(TraitProperty::device_isa___ANY, &Twenty)}) { + APInt Zero(64, 0), Four(64, 4); + for (TraitProperty Property : + {TraitProperty::user_condition_true, TraitProperty::device_isa___ANY}) { VariantMatchInfo ParallelFor, ParallelForParallel; - ParallelFor.addTrait(TraitProperty::implementation_vendor_llvm, "", Score); - ParallelFor.addTrait(Property, "sse"); + ParallelFor.addTrait(TraitProperty::implementation_vendor_llvm, "", &Four); + ParallelFor.addTrait(Property, "sse", &Zero); ParallelFor.addTrait(TraitProperty::construct_parallel_parallel, ""); ParallelFor.addTrait(TraitProperty::construct_for_for, ""); ParallelForParallel.addTrait(TraitProperty::implementation_vendor_llvm, ""); - ParallelForParallel.addTrait(Property, "sse2"); + ParallelForParallel.addTrait(Property, "sse2", &Zero); ParallelForParallel.addTrait(TraitProperty::construct_parallel_parallel, ""); ParallelForParallel.addTrait(TraitProperty::construct_for_for, ""); @@ -424,18 +658,16 @@ TEST_F(OpenMPContextTest, RepeatedConstructWithConditionOrISA) { } TEST_F(OpenMPContextTest, ConstructOrderInStrictSubset) { - APInt Five(64, 5), Nineteen(64, 19); + APInt Five(64, 5); { OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); - Context.addTrait(TraitProperty::construct_parallel_parallel); Context.addTrait(TraitProperty::construct_for_for); Context.addTrait(TraitProperty::construct_parallel_parallel); Context.addTrait(TraitProperty::construct_for_for); Context.addTrait(TraitProperty::construct_simd_simd); VariantMatchInfo ParallelFor, ForParallelSimd; - ParallelFor.addTrait(TraitProperty::implementation_vendor_llvm, "", - &Nineteen); + ParallelFor.addTrait(TraitProperty::implementation_vendor_llvm, "", &Five); ParallelFor.addTrait(TraitProperty::construct_parallel_parallel, ""); ParallelFor.addTrait(TraitProperty::construct_for_for, ""); ForParallelSimd.addTrait(TraitProperty::implementation_vendor_llvm, ""); @@ -444,7 +676,7 @@ TEST_F(OpenMPContextTest, ConstructOrderInStrictSubset) { ForParallelSimd.addTrait(TraitProperty::construct_parallel_parallel, ""); ForParallelSimd.addTrait(TraitProperty::construct_simd_simd, ""); - // Both score 23. Each property of ParallelFor is also required by + // Both score 12. Each property of ParallelFor is also required by // ForParallelSimd, but its construct sequence is not a subsequence. SmallVector<VariantMatchInfo, 2> Candidates{ParallelFor, ForParallelSimd}; EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); >From a3e97a634a41e7554aae1a899809977f3fd593f8 Mon Sep 17 00:00:00 2001 From: "Chi Chun, Chen" <[email protected]> Date: Wed, 21 Oct 2026 12:47:44 -0500 Subject: [PATCH 2/2] Test device weights at the TARGET boundary KIND's weight depends on the depth of the construct context, which starts at the innermost enclosing TARGET. A PARALLEL outside the TARGET must not raise it, and a PARALLEL inside must. --- .../OpenMP/declare_variant_construct_target.c | 38 +++++++++++++++++++ .../Lower/OpenMP/metadirective-construct.f90 | 35 +++++++++++++++++ 2 files changed, 73 insertions(+) diff --git a/clang/test/OpenMP/declare_variant_construct_target.c b/clang/test/OpenMP/declare_variant_construct_target.c index 1f3db3f8179cfc..dde4adb4239e1e 100644 --- a/clang/test/OpenMP/declare_variant_construct_target.c +++ b/clang/test/OpenMP/declare_variant_construct_target.c @@ -27,5 +27,43 @@ void inside(void) { base_inside(); } +// Device weights also count only from the innermost TARGET. KIND is worth 2 +// with only TARGET in the context and loses to the vendor score of 3. With +// PARALLEL inside TARGET, it is worth 4 and wins. + +#ifdef __NVPTX__ +#define KIND gpu +#else +#define KIND cpu +#endif + +int kind_outside(void); +int vendor_outside(void); +#pragma omp declare variant(kind_outside) match(device={kind(KIND)}) +#pragma omp declare variant(vendor_outside) \ + match(implementation={vendor(score(3): llvm)}) +int weight_outside(void); + +int kind_inside(void); +int vendor_inside(void); +#pragma omp declare variant(kind_inside) match(device={kind(KIND)}) +#pragma omp declare variant(vendor_inside) \ + match(implementation={vendor(score(3): llvm)}) +int weight_inside(void); + +void outside_weight(void) { +#pragma omp parallel +#pragma omp target + weight_outside(); +} + +void inside_weight(void) { +#pragma omp target +#pragma omp parallel + weight_inside(); +} + // CHECK-DAG: call{{.*}} @base_outside( // CHECK-DAG: call{{.*}} @parallel_inside( +// CHECK-DAG: call{{.*}} @vendor_outside( +// CHECK-DAG: call{{.*}} @kind_inside( diff --git a/flang/test/Lower/OpenMP/metadirective-construct.f90 b/flang/test/Lower/OpenMP/metadirective-construct.f90 index 66fab4bad0709d..2faae37580b8ba 100644 --- a/flang/test/Lower/OpenMP/metadirective-construct.f90 +++ b/flang/test/Lower/OpenMP/metadirective-construct.f90 @@ -172,6 +172,41 @@ subroutine test_construct_after_target() !$omp end parallel end subroutine +! Device weights also count only from the innermost TARGET. KIND is worth 2 here +! and loses to the vendor score of 3. +! CHECK-LABEL: func.func @_QPtest_device_weight_outside_target() +! CHECK: omp.parallel +! CHECK: omp.target +! CHECK-NOT: omp.taskyield +! CHECK: omp.taskwait +! CHECK: return +subroutine test_device_weight_outside_target() + !$omp parallel + !$omp target + !$omp metadirective & + !$omp & when(device={kind(host)}: taskyield) & + !$omp & when(implementation={vendor(score(3): llvm)}: taskwait) + !$omp end target + !$omp end parallel +end subroutine + +! With PARALLEL inside TARGET, KIND is worth 4 and wins. +! CHECK-LABEL: func.func @_QPtest_device_weight_inside_target() +! CHECK: omp.target +! CHECK: omp.parallel +! CHECK-NOT: omp.taskwait +! CHECK: omp.taskyield +! CHECK: return +subroutine test_device_weight_inside_target() + !$omp target + !$omp parallel + !$omp metadirective & + !$omp & when(device={kind(host)}: taskyield) & + !$omp & when(implementation={vendor(score(3): llvm)}: taskwait) + !$omp end parallel + !$omp end target +end subroutine + ! Both candidates score 8. The first selector is a strict subset of the second ! because the second repeats PARALLEL. ! CHECK-LABEL: func.func @_QPtest_repeated_construct_subset() _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
