https://github.com/chichunchen updated https://github.com/llvm/llvm-project/pull/224431
>From b4782b36c19749fb3f9df1ecec17d3812d91fb04 Mon Sep 17 00:00:00 2001 From: "Chi-Chun, Chen" <[email protected]> Date: Wed, 16 Sep 2026 16:48:23 -0500 Subject: [PATCH] [clang][flang][OpenMP] Fix context selector matching and scoring Incorrect context selector matching and scoring can select the wrong variant function or metadirective replacement. Compute device-selector weights from the enclosing construct-context depth. For example, inside a `parallel` region: | Selector | Before | After | | ------------------------ | ------ | ----- | | `device={kind(cpu)}` | 2 | 3 | | `construct={parallel}` | 2 | 2 | Previously, the incorrect tie allowed candidate order to determine the winner. The device selector now receives the higher score. Track enclosing constructs even when they have no corresponding selector property. Stop at and include the innermost `target`, or retain the full enclosing context when there is no `target`. Exclude `section` separators: counting the optional first `section` directive would change scores for equivalent code. Use the highest-scoring ordered construct match, count each selector's score once, and use arbitrary-width arithmetic to avoid overflow. Apply strict-subset zeroing before selecting the highest score: ```text compatible = candidates matching the context for each candidate in compatible: if candidate is a strict subset of another compatible candidate: score[candidate] = 0 else: score[candidate] = compute_score(candidate, context) select the first maximum-score candidate, or none if compatible is empty ``` For example, consider these two matching selectors: ```text A: implementation={vendor(score(100): llvm)} B: implementation={vendor(score(1): llvm)}, user={condition(1)} ``` A has an unadjusted score of 101, while B scores 2. However, A is a strict subset of B, so A's final score is zero and B wins. Subset checks must therefore apply even when the unadjusted scores differ. Preserve construct order and property identity in subset comparisons, and treat `kind(any)` as an absent kind selector. Retain unknown properties and their selector scores for `match_any` and `match_none`. Apply the same ranking rules to Flang's dynamic conditions and implicit `nothing`. Add shared-context unit tests and Clang C/C++ and Flang regression tests. This separates the matching and scoring fixes from the Flang lowering work in #219014. Assisted with Codex. --- clang/lib/AST/OpenMPClause.cpp | 25 +- clang/lib/Sema/SemaOpenMP.cpp | 47 +- .../declare_variant_device_kind_codegen.cpp | 5 +- clang/test/OpenMP/declare_variant_scoring.c | 206 ++++++++ clang/test/OpenMP/dispatch_variant_matching.c | 52 ++ flang/lib/Lower/OpenMP/Utils.cpp | 27 +- flang/lib/Lower/OpenMP/Utils.h | 4 +- flang/lib/Semantics/openmp-utils.cpp | 61 +-- .../OpenMP/declare-variant-construct.f90 | 169 ++++++- .../OpenMP/metadirective-device-kind.f90 | 17 + .../OpenMP/metadirective-implementation.f90 | 12 + .../OpenMP/metadirective-target-boundary.f90 | 62 +++ .../test/Lower/OpenMP/metadirective-user.f90 | 192 +++++++- .../OpenMP/variant-scoring-any-sections.f90 | 81 ++++ .../include/llvm/Frontend/OpenMP/OMPContext.h | 61 ++- llvm/lib/Frontend/OpenMP/OMPContext.cpp | 300 +++++++----- llvm/unittests/Frontend/OpenMPContextTest.cpp | 459 +++++++++++++++++- 17 files changed, 1584 insertions(+), 196 deletions(-) create mode 100644 clang/test/OpenMP/declare_variant_scoring.c create mode 100644 clang/test/OpenMP/dispatch_variant_matching.c create mode 100644 flang/test/Lower/OpenMP/metadirective-target-boundary.f90 create mode 100644 flang/test/Lower/OpenMP/variant-scoring-any-sections.f90 diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp index 2061d5395ac65..4658a77a40ee1 100644 --- a/clang/lib/AST/OpenMPClause.cpp +++ b/clang/lib/AST/OpenMPClause.cpp @@ -3302,8 +3302,29 @@ TargetOMPContext::TargetOMPContext( DiagUnknownTrait(std::move(DiagUnknownTrait)) { ASTCtx.getFunctionFeatureMap(FeatureMap, CurrentFunctionDecl); - for (llvm::omp::TraitProperty Property : ConstructTraits) - addTrait(Property); + // The construct context starts at and includes the innermost target: + // + // Enclosing stack: parallel -> target -> teams -> parallel + // Matching context: target -> teams -> parallel + // + // Constructs outside that target must not participate in matching or + // increase scoring depth. With no target, retain the entire stack. + auto Target = llvm::find(llvm::reverse(ConstructTraits), + llvm::omp::TraitProperty::construct_target_target); + if (Target != ConstructTraits.rend()) + ConstructTraits = ConstructTraits.take_back( + std::distance(ConstructTraits.rbegin(), Target) + 1); + + // Constructs without selector properties still occupy scoring positions. + // For example, parallel -> task has depth two, with task represented by + // an invalid placeholder. Record its position without activating invalid + // as a matchable trait. + for (llvm::omp::TraitProperty Property : ConstructTraits) { + if (Property == llvm::omp::TraitProperty::invalid) + addUnknownConstruct(); + else + addTrait(Property); + } } bool TargetOMPContext::matchesISATrait(StringRef RawString) const { diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 2e4d9f2f82f0b..8bac4a5ee5b6a 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -4448,16 +4448,43 @@ 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); + // Update the enclosing construct stack for declare variant matching and + // scoring on region entry or exit. Record each directive's constructs in + // nesting order, using placeholders for constructs without selector + // properties so they still contribute to scoring positions and depth. + for (OpenMPDirectiveKind Leaf : getLeafConstructsOrSelf(DKind)) { + switch (Leaf) { + case OMPD_target: + Traits.push_back(llvm::omp::TraitProperty::construct_target_target); + break; + case OMPD_teams: + Traits.push_back(llvm::omp::TraitProperty::construct_teams_teams); + break; + case OMPD_parallel: + Traits.push_back(llvm::omp::TraitProperty::construct_parallel_parallel); + break; + case OMPD_for: + Traits.push_back(llvm::omp::TraitProperty::construct_for_for); + break; + case OMPD_simd: + Traits.push_back(llvm::omp::TraitProperty::construct_simd_simd); + break; + case OMPD_section: + // SECTION separates blocks within SECTIONS and adds no construct level. + // Do not add a placeholder: spelling the optional first SECTION must + // not change variant scores. + break; + case OMPD_dispatch: + // OpenMP allows omitting DISPATCH from the construct context. Keep it + // omitted here: adding it for the whole region would also affect calls + // in arguments, but the trait may apply only to the target call. + break; + default: + // Constructs without a selector property still affect scoring depth. + Traits.push_back(llvm::omp::TraitProperty::invalid); + break; + } + } Stack->handleConstructTrait(Traits, ScopeEntry); } diff --git a/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp b/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp index 9335df10f957c..612a15002c3d6 100644 --- a/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp +++ b/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp @@ -108,8 +108,9 @@ #define WRONG host, nohost #endif // HOST #ifdef CPU -#define SUBSET cpu -#define CORRECT cpu, any +// kind(any) must appear alone and is equivalent to omitting kind. +#define SUBSET any +#define CORRECT cpu #define WRONG cpu, gpu #endif // CPU #ifdef NOHOST diff --git a/clang/test/OpenMP/declare_variant_scoring.c b/clang/test/OpenMP/declare_variant_scoring.c new file mode 100644 index 0000000000000..085e7b31e4b7f --- /dev/null +++ b/clang/test/OpenMP/declare_variant_scoring.c @@ -0,0 +1,206 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 \ +// RUN: -triple x86_64-unknown-linux -target-feature +avx \ +// RUN: -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -x c++ -verify -fopenmp -fopenmp-version=52 \ +// RUN: -triple x86_64-unknown-linux -target-feature +avx \ +// RUN: -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +#ifdef __cplusplus +extern "C" { +#endif + +#pragma omp begin declare target +void cpu_variant(void); +void scored_variant(void); +void parallel_variant(void); + +#pragma omp declare variant(cpu_variant) match(device = {kind(cpu)}) +#pragma omp declare variant(scored_variant) \ + match(implementation = {vendor(score(3) : llvm)}) +void target_base(void); + +#pragma omp declare variant(cpu_variant) match(device = {kind(cpu)}) +#pragma omp declare variant(parallel_variant) match(construct = {parallel}) +void depth_base(void); + +#pragma omp declare variant(parallel_variant) match(construct = {parallel}) +void construct_base(void); + +#pragma omp declare variant(cpu_variant) match(device = {kind(cpu)}) +#pragma omp declare variant(scored_variant) \ + match(implementation = {vendor(score(3) : llvm)}) +void task_depth_base(void); + +void subset_variant(void); +void superset_variant(void); +#pragma omp declare variant(subset_variant) \ + match(implementation = {vendor(score(100) : llvm)}) +#pragma omp declare variant(superset_variant) \ + match(implementation = {vendor(score(1) : llvm)}, user = {condition(1)}) +void subset_base(void); + +#pragma omp declare variant(subset_variant) \ + match(implementation = {vendor(score(100) : llvm)}) +#pragma omp declare variant(superset_variant) \ + match(implementation = {vendor(score(1) : llvm)}, device = {kind(any)}) +void any_base(void); + +#pragma omp declare variant(cpu_variant) match(device = {kind(cpu)}) +#pragma omp declare variant(scored_variant) \ + match(implementation = {vendor(score(5) : llvm)}) +void sections_base(void); + +void ordered_high_variant(void); +void ordered_low_variant(void); +#pragma omp declare variant(ordered_high_variant) \ + match(construct = {parallel, for}, \ + implementation = {vendor(score(100) : llvm)}) +#pragma omp declare variant(ordered_low_variant) \ + match(construct = {for, parallel, simd}, \ + implementation = {vendor(score(1) : llvm)}) +void ordered_base(void); + +void isa_high_variant(void); +void isa_low_variant(void); +#pragma omp declare variant(isa_high_variant) \ + match(device = {isa("sse2")}, \ + implementation = {vendor(score(100) : llvm)}) +#pragma omp declare variant(isa_low_variant) \ + match(device = {isa("avx")}, implementation = {vendor(score(1) : llvm)}, \ + user = {condition(1)}) +void isa_base(void); +#pragma omp end declare target + +// With only TARGET in the context, CPU scores 3 and the vendor variant's +// total score is 4. +void target_only(void) { +#pragma omp target + { target_base(); } +} +// CHECK-LABEL: define internal void @__omp_offloading_{{.*}}target_only +// CHECK: call void @scored_variant() +// CHECK: ret void + +// The outer PARALLEL must not increase the device score inside TARGET. +void parallel_target(void) { +#pragma omp parallel + { +#pragma omp target + { target_base(); } + } +} +// CHECK-LABEL: define internal void @__omp_offloading_{{.*}}parallel_target +// CHECK: call void @scored_variant() +// CHECK: ret void + +// In PARALLEL, CPU scores 3 and construct={parallel} scores 2. +void parallel_depth(void) { +#pragma omp parallel + { depth_base(); } +} +// CHECK-LABEL: define internal void @parallel_depth.omp_outlined +// CHECK: call void @cpu_variant() +// CHECK: ret void + +// TARGET is retained, as are constructs nested inside it: CPU scores 5. +void target_parallel(void) { +#pragma omp target + { +#pragma omp parallel + { target_base(); } + } +} +// CHECK-LABEL: define internal void @__omp_offloading_{{.*}}target_parallel +// CHECK: define internal void @{{.*}}omp_outlined +// CHECK: call void @cpu_variant() +// CHECK: ret void + +// Leaving TARGET restores the enclosing PARALLEL context. +void restore_context(void) { +#pragma omp parallel + { +#pragma omp target + { target_base(); } + depth_base(); + construct_base(); + } +} +// CHECK-LABEL: define internal void @restore_context.omp_outlined +// CHECK: call void @cpu_variant() +// CHECK: call void @parallel_variant() +// CHECK: ret void +// CHECK-LABEL: define internal void @__omp_offloading_{{.*}}restore_context +// CHECK: call void @scored_variant() +// CHECK: ret void + +// TASK has no construct-selector property, but contributes to the device +// weight. In PARALLEL > TASK, CPU scores 5 and beats the vendor variant's +// total score of 4. +void task_depth(void) { +#pragma omp parallel + { +#pragma omp task + { task_depth_base(); } + } +} +// CHECK-LABEL: define internal {{.*}}i32 @.omp_task_entry. +// CHECK: call void @cpu_variant() + +// A strict subset has score zero before candidates are ranked, even when its +// explicit score would otherwise be higher. +void strict_subset(void) { subset_base(); } +// CHECK-LABEL: define{{.*}} void @strict_subset +// CHECK: call void @superset_variant() +// CHECK: ret void + +// kind(any) does not make the lower-scored selector a strict superset. +void kind_any(void) { any_base(); } +// CHECK-LABEL: define{{.*}} void @kind_any +// CHECK: call void @subset_variant() +// CHECK: ret void + +// SECTION is a separator, so CPU scores 5 and the vendor scores 6 in both +// spellings of the first section of PARALLEL SECTIONS. +void explicit_section(void) { +#pragma omp parallel sections + { +#pragma omp section + { sections_base(); } + } +} +// CHECK-LABEL: define internal void @explicit_section.omp_outlined +// CHECK: call void @scored_variant() +// CHECK: ret void + +void implicit_section(void) { +#pragma omp parallel sections + { sections_base(); } +} +// CHECK-LABEL: define internal void @implicit_section.omp_outlined +// CHECK: call void @scored_variant() +// CHECK: ret void + +// Different construct orders do not create a subset relationship. +void different_construct_order(void) { +#pragma omp parallel for + for (int i = 0; i < 2; ++i) { +#pragma omp parallel for simd + for (int j = 0; j < 2; ++j) + ordered_base(); + } +} +// CHECK-LABEL: define internal void @different_construct_order.omp_outlined +// CHECK: define internal void @{{.*}}omp_outlined +// CHECK: call void @ordered_high_variant() +// CHECK: ret void + +// Different active ISA properties do not create a subset relationship. +void different_isa(void) { isa_base(); } +// CHECK-LABEL: define{{.*}} void @different_isa +// CHECK: call void @isa_high_variant() +// CHECK: ret void + +#ifdef __cplusplus +} +#endif diff --git a/clang/test/OpenMP/dispatch_variant_matching.c b/clang/test/OpenMP/dispatch_variant_matching.c new file mode 100644 index 0000000000000..fd2664be2ae79 --- /dev/null +++ b/clang/test/OpenMP/dispatch_variant_matching.c @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -verify -ast-dump %s \ +// RUN: | FileCheck %s --implicit-check-not=PseudoObjectExpr +// RUN: %clang_cc1 -x c++ -fopenmp -fopenmp-version=52 -verify -ast-dump %s \ +// RUN: | FileCheck %s --implicit-check-not=PseudoObjectExpr +// expected-no-diagnostics + +// Clang omits the implementation-defined DISPATCH construct trait. In +// particular, it must not select dispatch variants in argument expressions. +int g_variant(void); +#pragma omp declare variant(g_variant) match(construct = {dispatch}) +int g(void); + +int f_variant(int); +#pragma omp declare variant(f_variant) match(construct = {dispatch}) +int f(int); +void plain(int); + +// CHECK-LABEL: FunctionDecl {{.*}} test_argument +// CHECK: OMPDispatchDirective +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'plain' +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'g' +void test_argument(void) { +#pragma omp dispatch + plain(g()); +} + +// CHECK-LABEL: FunctionDecl {{.*}} test_target +// CHECK: OMPDispatchDirective +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'f' +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'g' +void test_target(void) { +#pragma omp dispatch + f(g()); +} + +// CHECK-LABEL: FunctionDecl {{.*}} test_assignment +// CHECK: OMPDispatchDirective +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'f' +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'g' +void test_assignment(int *result) { +#pragma omp dispatch + *result = f(g()); +} + +// CHECK-LABEL: FunctionDecl {{.*}} test_cast +// CHECK: OMPDispatchDirective +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'f' +// CHECK: DeclRefExpr {{.*}} Function {{.*}} 'g' +void test_cast(void) { +#pragma omp dispatch + (void)f(g()); +} diff --git a/flang/lib/Lower/OpenMP/Utils.cpp b/flang/lib/Lower/OpenMP/Utils.cpp index b52c03ff63fed..ca02a069a3f1b 100644 --- a/flang/lib/Lower/OpenMP/Utils.cpp +++ b/flang/lib/Lower/OpenMP/Utils.cpp @@ -1481,17 +1481,36 @@ void collectEnclosingConstructTraits( // be able to match construct={target, parallel}. The final reverse yields // outermost-to-innermost order as required by OMPContext. for (; op; op = op->getParentOp()) { - if (mlir::isa<mlir::omp::WsloopOp>(op)) + if (mlir::isa<mlir::omp::SimdOp>(op)) + constructTraits.push_back(llvm::omp::TraitProperty::construct_simd_simd); + else if (mlir::isa<mlir::omp::WsloopOp>(op)) constructTraits.push_back(llvm::omp::TraitProperty::construct_for_for); - if (mlir::isa<mlir::omp::ParallelOp>(op)) + else if (mlir::isa<mlir::omp::ParallelOp>(op)) constructTraits.push_back( llvm::omp::TraitProperty::construct_parallel_parallel); - if (mlir::isa<mlir::omp::TeamsOp>(op)) + else if (mlir::isa<mlir::omp::TeamsOp>(op)) constructTraits.push_back( llvm::omp::TraitProperty::construct_teams_teams); - if (mlir::isa<mlir::omp::TargetOp>(op)) + else if (mlir::isa<mlir::omp::TargetOp>(op)) { constructTraits.push_back( llvm::omp::TraitProperty::construct_target_target); + // The construct context starts at the innermost TARGET, as in + // semantic analysis. + break; + } else if (mlir::isa<mlir::omp::CriticalOp, mlir::omp::DistributeOp, + mlir::omp::FuseOp, mlir::omp::LoopOp, + mlir::omp::MaskedOp, mlir::omp::MasterOp, + mlir::omp::OrderedRegionOp, mlir::omp::ScopeOp, + mlir::omp::SectionsOp, mlir::omp::SingleOp, + mlir::omp::TargetDataOp, mlir::omp::TaskgroupOp, + mlir::omp::TaskloopContextOp, mlir::omp::TaskOp, + mlir::omp::TileOp, mlir::omp::UnrollFullOp, + mlir::omp::UnrollPartialOp, + mlir::omp::WorkdistributeOp, mlir::omp::WorkshareOp>( + op)) + // These source constructs have no construct-selector property, but + // still occupy a position and contribute to device-selector weights. + constructTraits.push_back(llvm::omp::TraitProperty::invalid); } std::reverse(constructTraits.begin(), constructTraits.end()); } diff --git a/flang/lib/Lower/OpenMP/Utils.h b/flang/lib/Lower/OpenMP/Utils.h index 7fde1825b006a..f630691421c34 100644 --- a/flang/lib/Lower/OpenMP/Utils.h +++ b/flang/lib/Lower/OpenMP/Utils.h @@ -265,7 +265,9 @@ std::optional<llvm::SmallVector<mlir::Value>> getIteratorElementIndices( /// Walk the already-emitted MLIR parent operations starting from \p op and /// collect the implied OpenMP construct traits in outermost-to-innermost /// order. Used by metadirective lowering and declare-variant call resolution -/// to build the `ConstructTraits` of an `OMPContext`. +/// to build the `ConstructTraits` of an `OMPContext`. Constructs without a +/// corresponding trait property are represented by `TraitProperty::invalid` +/// so they still contribute to scoring positions and depth. void collectEnclosingConstructTraits( mlir::Operation *op, llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits); diff --git a/flang/lib/Semantics/openmp-utils.cpp b/flang/lib/Semantics/openmp-utils.cpp index a341cf4fa1706..136124664f4b0 100644 --- a/flang/lib/Semantics/openmp-utils.cpp +++ b/flang/lib/Semantics/openmp-utils.cpp @@ -2389,10 +2389,9 @@ void ProcessTraitProperties(llvm::omp::VariantMatchInfo &vmi, vmi.addTrait(set, llvm::omp::TraitProperty::target_device_isa___ANY, name->v, scorePtr); } else { - // For non-ISA selectors (arch, kind, vendor, etc.), unknown properties - // mean the variant cannot match. Add an invalid trait to ensure it is - // not selected. - vmi.addTrait(llvm::omp::TraitProperty::invalid, name->v, scorePtr); + // Unknown properties remain inactive, but their selectors still + // contribute to scoring under match_any or match_none. + vmi.addUnknownTrait(selector, name->v, scorePtr); } } } @@ -2480,17 +2479,18 @@ static void AddTraitPropertiesFromSelector(llvm::omp::TraitSet set, continue; } if (auto constValue{EvaluateUserCondition(semaCtx, *scalarExpr)}) { + llvm::StringRef condition{prop.source.begin(), prop.source.size()}; vmi.addTrait(set, *constValue ? llvm::omp::TraitProperty::user_condition_true : llvm::omp::TraitProperty::user_condition_false, - "<condition>", scorePtr); + condition, scorePtr); continue; } if (!dynamicCond) { dynamicCond = DynamicUserCondition{scalarExpr, prop.source}; } vmi.addTrait(set, llvm::omp::TraitProperty::user_condition_unknown, - "<condition>", scorePtr); + llvm::StringRef{prop.source.begin(), prop.source.size()}, scorePtr); } return; } @@ -2505,6 +2505,13 @@ static void AddTraitPropertiesFromSelector(llvm::omp::TraitSet set, // the selector itself implies the property. if (const auto *dir{std::get_if<llvm::omp::Directive>(&traitName.u)}) { AppendConstructTraitsForDirective(*dir, vmi); + } else if (const auto *value{std::get_if<parser::OmpTraitSelectorName::Value>( + &traitName.u)}) { + // SIMD is a predefined selector name because it can take clause + // properties, unlike the other construct selectors. + if (*value == parser::OmpTraitSelectorName::Value::Simd) { + AppendConstructTraitsForDirective(llvm::omp::Directive::OMPD_simd, vmi); + } } } @@ -2600,8 +2607,11 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet( staticVMI.ScoreMap.erase(scoreIt); } staticVMI.RequiredTraits.reset(unsigned(dynamicConditionTrait)); + staticVMI.UserCondition = {}; llvm::APInt *conditionScorePtr{ conditionScore ? &*conditionScore : nullptr}; + llvm::StringRef conditionSource{ + dynamicCondition->source.begin(), dynamicCondition->source.size()}; bool hasMatchAny{rawVMI.RequiredTraits.test(unsigned(matchAnyTrait))}; bool hasMatchNone{rawVMI.RequiredTraits.test(unsigned(matchNoneTrait))}; @@ -2610,15 +2620,13 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet( // Only match_any can remain applicable when the static traits do not // match, because a true runtime condition may satisfy the selector. if (!isStaticVMIApplicable) { - if (!hasMatchAny || - staticVMI.RequiredTraits.test( - unsigned(llvm::omp::TraitProperty::invalid))) { + if (!hasMatchAny) { continue; } llvm::omp::VariantMatchInfo conditionTrueVMI{staticVMI}; conditionTrueVMI.addTrait( - llvm::omp::TraitProperty::user_condition_true, "<condition>", + llvm::omp::TraitProperty::user_condition_true, conditionSource, conditionScorePtr); if (!llvm::omp::isVariantApplicableInContext( conditionTrueVMI, matchContext)) { @@ -2631,31 +2639,24 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet( rankingVMI.addTrait(hasMatchNone ? dynamicConditionTrait : llvm::omp::TraitProperty::user_condition_true, - "<condition>", conditionScorePtr); + conditionSource, conditionScorePtr); }; if (hasMatchAny && isStaticVMIApplicable) { // Represent both outcomes: a guarded candidate with the condition's - // score and an unguarded candidate with only the static traits. If - // the WHEN clause omits its directive, only add the unguarded - // candidate. - if (isExplicit) { - llvm::omp::VariantMatchInfo conditionTrueVMI{staticVMI}; - addConditionTraitForRanking(conditionTrueVMI); - result.candidates.push_back({spec, std::move(conditionTrueVMI), - isExplicit, dynamicCondition}); - } + // score and an unguarded candidate with only the static traits. + llvm::omp::VariantMatchInfo conditionTrueVMI{staticVMI}; + addConditionTraitForRanking(conditionTrueVMI); + result.candidates.push_back({spec, std::move(conditionTrueVMI), + isExplicit, dynamicCondition}); result.candidates.push_back({spec, std::move(staticVMI), isExplicit}); continue; } llvm::omp::VariantMatchInfo rankingVMI{staticVMI}; - // Preserve the existing lowering behavior for an omitted directive: - // do not let its runtime condition raise the implicit NOTHING rank. - if (!isExplicit && hasMatchAny && !isStaticVMIApplicable) - rankingVMI = llvm::omp::VariantMatchInfo(); - else if (isExplicit) - addConditionTraitForRanking(rankingVMI); + // Implicit NOTHING participates in scoring just like an explicit + // replacement; explicitness only breaks ties between equal scores. + addConditionTraitForRanking(rankingVMI); result.candidates.push_back({spec, std::move(rankingVMI), isExplicit, dynamicCondition, /*conditionShouldBeTrue=*/!hasMatchNone}); continue; @@ -2889,7 +2890,8 @@ bool MayVariantBeSelected( bool userTrue{required.test(unsigned(TP::user_condition_true))}; bool userUnknown{required.test(unsigned(TP::user_condition_unknown))}; bool userFalse{required.test(unsigned(TP::user_condition_false))}; - bool invalid{required.test(unsigned(TP::invalid))}; + bool invalid{ + required.test(unsigned(TP::invalid)) || !vmi.UnknownTraits.empty()}; // The target-only LLVM matcher below skips user and construct traits while // retaining the global match kind. Account for those skipped traits first; @@ -2949,7 +2951,10 @@ OmpVariantMatchContext::OmpVariantMatchContext(bool isDeviceCompilation, std::move(targetOffloadTriple), /*DeviceNum=*/-1), features_(std::move(targetFeatures)) { for (llvm::omp::TraitProperty trait : constructTraits) { - addTrait(trait); + if (trait == llvm::omp::TraitProperty::invalid) + addUnknownConstruct(); + else + addTrait(trait); } } diff --git a/flang/test/Lower/OpenMP/declare-variant-construct.f90 b/flang/test/Lower/OpenMP/declare-variant-construct.f90 index ca72b5c8e47af..90fe7250b9bd6 100644 --- a/flang/test/Lower/OpenMP/declare-variant-construct.f90 +++ b/flang/test/Lower/OpenMP/declare-variant-construct.f90 @@ -1,8 +1,8 @@ ! RUN: %flang_fc1 -emit-fir -fopenmp -fopenmp-version=51 %s -o - | FileCheck %s ! RUN: %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=51 %s -o - | FileCheck %s -! DECLARE VARIANT callee resolution with combined/composite construct -! selectors. The bases and their variants are sibling module procedures, so +! DECLARE VARIANT callee resolution with construct selectors, including +! combined/composite selectors. Bases and variants are sibling procedures, so ! each variant is accessible at every reference to its base. module m @@ -10,6 +10,79 @@ module m subroutine base_tt !$omp declare variant (base_tt:vsub_tt) match (construct={target teams}) end subroutine base_tt + + subroutine base_device_weight + !$omp declare variant (vsub_cpu) match(device={kind(cpu)}) + !$omp declare variant (vsub_parallel) match(construct={parallel}) + end subroutine + subroutine vsub_cpu + end subroutine + subroutine vsub_parallel + end subroutine + + ! CPU scores 5 in this context, beating PARALLEL's score of 3. + ! CHECK-LABEL: func.func @_QMmPtest_device_weight() + ! CHECK: omp.parallel + ! CHECK: omp.parallel + ! CHECK-NOT: fir.call @_QMmPvsub_parallel + ! CHECK: fir.call @_QMmPvsub_cpu() + ! CHECK-NOT: fir.call @_QMmPvsub_parallel + ! CHECK: return + subroutine test_device_weight + !$omp parallel + !$omp parallel + call base_device_weight() + !$omp end parallel + !$omp end parallel + end subroutine + + subroutine base_task_weight + !$omp declare variant (vsub_cpu) match(device={kind(cpu)}) + !$omp declare variant (vsub_score3) & + !$omp& match(implementation={vendor(score(3): llvm)}) + end subroutine + subroutine vsub_score3 + end subroutine + + ! TASK has no construct-selector property, but contributes to the device + ! weight. In PARALLEL > TASK, CPU scores 5 and beats score(3)'s total of 4. + ! CHECK-LABEL: func.func @_QMmPtest_task_weight() + ! CHECK: omp.parallel + ! CHECK: omp.task + ! CHECK-NOT: fir.call @_QMmPvsub_score3 + ! CHECK: fir.call @_QMmPvsub_cpu() + ! CHECK-NOT: fir.call @_QMmPvsub_score3 + ! CHECK: return + subroutine test_task_weight + !$omp parallel + !$omp task + call base_task_weight() + !$omp end task + !$omp end parallel + end subroutine + + subroutine base_subset + !$omp declare variant (vsub_subset) & + !$omp& match(implementation={vendor(score(100): llvm)}) + !$omp declare variant (vsub_superset) & + !$omp& match(implementation={vendor(score(1): llvm)}, & + !$omp& user={condition(.true.)}) + end subroutine + subroutine vsub_subset + end subroutine + subroutine vsub_superset + end subroutine + + ! A strict subset has score zero before candidates are ranked, even when its + ! explicit score would otherwise be higher. + ! CHECK-LABEL: func.func @_QMmPtest_strict_subset() + ! CHECK-NOT: fir.call @_QMmPvsub_subset + ! CHECK: fir.call @_QMmPvsub_superset() + ! CHECK-NOT: fir.call @_QMmPvsub_subset + subroutine test_strict_subset + call base_subset() + end subroutine + subroutine vsub_tt end subroutine vsub_tt @@ -49,6 +122,38 @@ end subroutine vsub_lo subroutine vsub_hi end subroutine vsub_hi + subroutine base_simd + !$omp declare variant (vsub_simd) match (construct={simd}) + end subroutine base_simd + subroutine vsub_simd + end subroutine vsub_simd + + subroutine base_do_simd + !$omp declare variant (vsub_do_simd) match (construct={do, simd}) + end subroutine base_do_simd + subroutine vsub_do_simd + end subroutine vsub_do_simd + + subroutine base_repeated + !$omp declare variant (vsub_lo) & + !$omp& match(user={condition(score(1): .true.)}) + !$omp declare variant (vsub_par) match(construct={parallel}) + end subroutine base_repeated + + ! The inner PARALLEL raises the construct score above the user score. + ! CHECK-LABEL: func.func @_QMmPtest_repeated_parallel( + ! CHECK: omp.parallel + ! CHECK: omp.parallel + ! CHECK-NOT: fir.call @_QMmPvsub_lo + ! CHECK: fir.call @_QMmPvsub_par() + subroutine test_repeated_parallel + !$omp parallel + !$omp parallel + call base_repeated() + !$omp end parallel + !$omp end parallel + end subroutine test_repeated_parallel + ! The combined directive selector decomposes to {target, teams}; it matches ! only when both constructs enclose the call. @@ -157,4 +262,64 @@ end subroutine test_rank_parallel_only subroutine test_score_ranking call base_score() end subroutine test_score_ranking + + ! Without an enclosing SIMD construct, neither selector matches. + + ! CHECK-LABEL: func.func @_QMmPtest_outside_simd( + ! CHECK: fir.call @_QMmPbase_simd() + ! CHECK: fir.call @_QMmPbase_do_simd() + subroutine test_outside_simd + call base_simd() + call base_do_simd() + end subroutine test_outside_simd + + ! SIMD supplies its construct trait, but not the DO trait. + + ! CHECK-LABEL: func.func @_QMmPtest_inside_simd( + ! CHECK: omp.simd + ! CHECK: omp.loop_nest + ! CHECK: fir.call @_QMmPvsub_simd() + ! CHECK: fir.call @_QMmPbase_do_simd() + subroutine test_inside_simd(n) + integer :: n, i + !$omp simd + do i = 1, n + call base_simd() + call base_do_simd() + end do + end subroutine test_inside_simd + + ! DO SIMD supplies both traits in DO -> SIMD order. + + ! CHECK-LABEL: func.func @_QMmPtest_inside_do_simd( + ! CHECK: omp.wsloop + ! CHECK: omp.simd + ! CHECK: omp.loop_nest + ! CHECK: fir.call @_QMmPvsub_simd() + ! CHECK: fir.call @_QMmPvsub_do_simd() + subroutine test_inside_do_simd(n) + integer :: n, i + !$omp do simd + do i = 1, n + call base_simd() + call base_do_simd() + end do + end subroutine test_inside_do_simd + + ! TARGET hides the outer PARALLEL during callee selection as well. + + ! CHECK-LABEL: func.func @_QMmPtest_target_boundary( + ! CHECK: omp.parallel + ! CHECK: omp.target + ! CHECK: fir.call @_QMmPbase_tp2() + ! CHECK-NEXT: fir.call @_QMmPbase_rank() + ! CHECK: return + subroutine test_target_boundary + !$omp parallel + !$omp target + call base_tp2() + call base_rank() + !$omp end target + !$omp end parallel + end subroutine test_target_boundary end module m diff --git a/flang/test/Lower/OpenMP/metadirective-device-kind.f90 b/flang/test/Lower/OpenMP/metadirective-device-kind.f90 index ae7613d40c5d9..d96be7739ac98 100644 --- a/flang/test/Lower/OpenMP/metadirective-device-kind.f90 +++ b/flang/test/Lower/OpenMP/metadirective-device-kind.f90 @@ -1,5 +1,22 @@ ! RUN: %flang_fc1 -fopenmp -emit-hlfir -fopenmp-version=50 %s -o - | FileCheck %s +! Device weights depend on context depth, not the candidate's selector count. +! CHECK-LABEL: func.func @_QPtest_device_kind_nested_parallel() +! CHECK: omp.parallel +! CHECK: omp.parallel +! CHECK-NOT: omp.barrier +! CHECK: omp.taskyield +! CHECK-NOT: omp.barrier +! CHECK: return +subroutine test_device_kind_nested_parallel() + !$omp parallel + !$omp parallel + !$omp metadirective when(device={kind(cpu)}: taskyield) & + !$omp& when(construct={parallel}: barrier) + !$omp end parallel + !$omp end parallel +end subroutine + ! CHECK-LABEL: func.func @_QPtest_device_kind_host() ! CHECK: omp.taskyield ! CHECK: return diff --git a/flang/test/Lower/OpenMP/metadirective-implementation.f90 b/flang/test/Lower/OpenMP/metadirective-implementation.f90 index af8f2af938ca5..cd61b733cbb1d 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/flang/test/Lower/OpenMP/metadirective-target-boundary.f90 b/flang/test/Lower/OpenMP/metadirective-target-boundary.f90 new file mode 100644 index 0000000000000..d7963bf7fe4ec --- /dev/null +++ b/flang/test/Lower/OpenMP/metadirective-target-boundary.f90 @@ -0,0 +1,62 @@ +! RUN: %flang_fc1 -fopenmp -fopenmp-version=51 -emit-fir %s -o - | FileCheck %s +! RUN: %flang_fc1 -fopenmp -fopenmp-version=51 -emit-hlfir %s -o - | \ +! RUN: FileCheck %s + +! TARGET hides the outer PARALLEL, so the SIMD replacement is not lowered. +! CHECK-LABEL: func.func @_QPactual_target( +! CHECK: omp.parallel +! CHECK: omp.target +! CHECK-NOT: omp.simd +! CHECK: return +subroutine actual_target(n, a) + integer :: n, i, a(n) + !$omp parallel + !$omp target + !$omp metadirective & + !$omp& when(construct={parallel}: simd) default(nothing) + do i = 1, n + a(i) = i + end do + !$omp end target + !$omp end parallel +end subroutine + +! A TARGET selected by a metadirective creates the same context boundary. +! CHECK-LABEL: func.func @_QPselected_target( +! CHECK: omp.parallel +! CHECK: omp.target +! CHECK-NOT: omp.simd +! CHECK: return +subroutine selected_target(n, a) + integer :: n, i, a(n) + !$omp parallel + !$omp begin metadirective default(target) + !$omp metadirective & + !$omp& when(construct={parallel, target}: simd) & + !$omp& default(nothing) + do i = 1, n + a(i) = i + end do + !$omp end metadirective + !$omp end parallel +end subroutine + +! The boundary includes TARGET itself and constructs nested inside it. +! CHECK-LABEL: func.func @_QPtarget_inner_parallel() +! CHECK: omp.parallel +! CHECK: omp.target +! CHECK: omp.parallel +! CHECK-NOT: omp.taskyield +! CHECK: omp.barrier +! CHECK-NOT: omp.taskyield +! CHECK: return +subroutine target_inner_parallel() + !$omp parallel + !$omp target + !$omp parallel + !$omp metadirective when(construct={target, parallel}: barrier) & + !$omp& default(taskyield) + !$omp end parallel + !$omp end target + !$omp end parallel +end subroutine diff --git a/flang/test/Lower/OpenMP/metadirective-user.f90 b/flang/test/Lower/OpenMP/metadirective-user.f90 index cdfbddd4151a0..060f5abc9587c 100644 --- a/flang/test/Lower/OpenMP/metadirective-user.f90 +++ b/flang/test/Lower/OpenMP/metadirective-user.f90 @@ -4,6 +4,185 @@ ! 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 +!===----------------------------------------------------------------------===! +! Unknown ARCH retains its weight even when only the runtime condition matches. +! CHECK-LABEL: func.func @_QPtest_unknown_arch_weight( +! CHECK: fir.if +! CHECK-NEXT: omp.barrier +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.taskyield +! CHECK: return +subroutine test_unknown_arch_weight(flag) + logical :: flag + !$omp metadirective & + !$omp& when(device={arch(bogus_arch)}, & + !$omp& implementation={extension(match_any)}, & + !$omp& user={condition(flag)}: barrier) & + !$omp& when(user={condition(score(1): .true.)}: taskyield) +end subroutine + +! Selectors with unknown properties retain their separate scores, +! regardless of selector order. +! CHECK-LABEL: func.func @_QPtest_unknown_selector_scores( +! CHECK: fir.if +! CHECK-NEXT: omp.barrier +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.taskyield +! CHECK: return +subroutine test_unknown_selector_scores(flag) + logical :: flag + !$omp metadirective & + !$omp& when(implementation={vendor(score(10): bogus_vendor), & + !$omp& extension(score(1): match_any, bogus_extension)}, & + !$omp& user={condition(score(5): flag)}: barrier) & + !$omp& when(implementation={vendor(score(10): llvm)}: taskyield) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_unknown_selector_scores_reversed( +! CHECK: fir.if +! CHECK-NEXT: omp.barrier +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.taskyield +! CHECK: return +subroutine test_unknown_selector_scores_reversed(flag) + logical :: flag + !$omp metadirective & + !$omp& when(implementation={extension(score(1): match_any, bogus_extension), & + !$omp& vendor(score(10): bogus_vendor)}, & + !$omp& user={condition(score(5): flag)}: barrier) & + !$omp& when(implementation={vendor(score(10): llvm)}: taskyield) +end subroutine + +! An unknown vendor does not veto a runtime MATCH_ANY condition or its score. +! CHECK-LABEL: func.func @_QPtest_dynamic_unknown_vendor( +! CHECK: fir.if +! CHECK-NEXT: omp.barrier +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.taskyield +! CHECK: return +subroutine test_dynamic_unknown_vendor(flag) + logical :: flag + !$omp metadirective & + !$omp& when(implementation={vendor(bogus_vendor), extension(match_any)}, & + !$omp& user={condition(score(5): flag)}: barrier) & + !$omp& when(user={condition(.true.)}: taskyield) +end subroutine + +! The same rule applies to unknown device traits and implicit NOTHING. +! CHECK-LABEL: func.func @_QPtest_dynamic_unknown_arch_implicit( +! CHECK: fir.if +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.barrier +! CHECK: return +subroutine test_dynamic_unknown_arch_implicit(flag) + logical :: flag + !$omp metadirective & + !$omp& when(device={arch(bogus_arch)}, & + !$omp& implementation={extension(match_any)}, & + !$omp& user={condition(score(10): flag)}:) & + !$omp& when(user={condition(score(5): .true.)}: barrier) +end subroutine + +! Large non-negative scores must not wrap their sum to zero. +! CHECK-LABEL: func.func @_QPtest_wide_score() +! CHECK: omp.parallel +! CHECK-NOT: omp.taskyield +! CHECK: omp.barrier +! CHECK-NOT: omp.taskyield +! CHECK: return +subroutine test_wide_score() + !$omp parallel + !$omp metadirective & + !$omp& when(user={condition(score(9223372036854775807_8): .true.)}, & + !$omp& implementation={vendor(score(9223372036854775807_8): llvm)}, & + !$omp& construct={parallel}: barrier) & + !$omp& when(user={condition(.true.)}: taskyield) + !$omp end parallel +end subroutine + +! Scored implicit NOTHING competes with explicit replacements by score. +!===----------------------------------------------------------------------===! + +! CHECK-LABEL: func.func @_QPtest_implicit_nothing_score( +! CHECK: fir.if +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.barrier +! CHECK: return +subroutine test_implicit_nothing_score(flag) + logical :: flag + !$omp metadirective & + !$omp& when(user={condition(score(10): flag)}:) & + !$omp& when(user={condition(score(5): .true.)}: barrier) +end subroutine + +! Explicit NOTHING with the same score must produce the same selection. +! CHECK-LABEL: func.func @_QPtest_explicit_nothing_score( +! CHECK: fir.if +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.barrier +! CHECK: return +subroutine test_explicit_nothing_score(flag) + logical :: flag + !$omp metadirective & + !$omp& when(user={condition(score(10): flag)}: nothing) & + !$omp& when(user={condition(score(5): .true.)}: barrier) +end subroutine + +! Equal scores favor the explicit replacement without a runtime branch. +! CHECK-LABEL: func.func @_QPtest_implicit_nothing_equal_score( +! CHECK-NOT: fir.if +! CHECK: omp.barrier +! CHECK-NEXT: return +subroutine test_implicit_nothing_equal_score(flag) + logical :: flag + !$omp metadirective & + !$omp& when(user={condition(score(5): flag)}:) & + !$omp& when(user={condition(score(5): .true.)}: barrier) +end subroutine + +! MATCH_ANY still needs a scored runtime candidate when a static trait matches. +! CHECK-LABEL: func.func @_QPtest_implicit_nothing_match_any_static( +! CHECK: fir.if +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.barrier +! CHECK: return +subroutine test_implicit_nothing_match_any_static(flag) + logical :: flag + !$omp metadirective & + !$omp& when(implementation={vendor(llvm), extension(match_any)}, & + !$omp& user={condition(score(10): flag)}:) & + !$omp& when(user={condition(score(5): .true.)}: barrier) +end subroutine + +! MATCH_ANY can also depend entirely on the runtime condition. +! CHECK-LABEL: func.func @_QPtest_implicit_nothing_match_any_runtime( +! CHECK: fir.if +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.barrier +! CHECK: return +subroutine test_implicit_nothing_match_any_runtime(flag) + logical :: flag + !$omp metadirective & + !$omp& when(implementation={vendor(gnu), extension(match_any)}, & + !$omp& user={condition(score(10): flag)}:) & + !$omp& when(user={condition(score(5): .true.)}: barrier) +end subroutine + +! MATCH_NONE retains the score but selects NOTHING when the condition is false. +! CHECK-LABEL: func.func @_QPtest_implicit_nothing_match_none( +! CHECK: arith.xori +! CHECK: fir.if +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.barrier +! CHECK: return +subroutine test_implicit_nothing_match_none(flag) + logical :: flag + !$omp metadirective & + !$omp& when(implementation={extension(match_none)}, & + !$omp& user={condition(score(10): flag)}:) & + !$omp& when(user={condition(score(5): .true.)}: barrier) +end subroutine + !===----------------------------------------------------------------------===! ! Static (constant-folded) user conditions !===----------------------------------------------------------------------===! @@ -425,13 +604,14 @@ subroutine test_dynamic_user_match_any_static_score(flag) #endif end subroutine -! The explicit directive variant wins this tie over the earlier implicit -! nothing candidate. -! CHECK-LABEL: func.func @_QPtest_dynamic_implicit_nothing_tie_break( -! CHECK-NOT: fir.if -! CHECK: omp.barrier +! The vendor-only selector is a strict subset of the implicit NOTHING's +! selector. The user condition determines which replacement is selected. +! CHECK-LABEL: func.func @_QPtest_dynamic_implicit_nothing_more_specific( +! CHECK: fir.if +! CHECK-NEXT: } else { +! CHECK-NEXT: omp.barrier ! CHECK: return -subroutine test_dynamic_implicit_nothing_tie_break(flag) +subroutine test_dynamic_implicit_nothing_more_specific(flag) logical, intent(in) :: flag !$omp metadirective & !$omp & when(implementation={vendor(llvm)}, user={condition(flag)}:) & diff --git a/flang/test/Lower/OpenMP/variant-scoring-any-sections.f90 b/flang/test/Lower/OpenMP/variant-scoring-any-sections.f90 new file mode 100644 index 0000000000000..e51df9bc37746 --- /dev/null +++ b/flang/test/Lower/OpenMP/variant-scoring-any-sections.f90 @@ -0,0 +1,81 @@ +! RUN: %flang_fc1 -emit-fir -fopenmp -fopenmp-version=52 %s -o - \ +! RUN: | FileCheck %s +! RUN: %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=52 %s -o - \ +! RUN: | FileCheck %s + +module scoring +contains + subroutine high() + end subroutine + subroutine low() + end subroutine + subroutine cpu() + end subroutine + subroutine vendor() + end subroutine + + subroutine base_any() + !$omp declare variant(high) & + !$omp& match(implementation={vendor(score(100): llvm)}) + !$omp declare variant(low) & + !$omp& match(implementation={vendor(score(1): llvm)}, device={kind(any)}) + end subroutine + + ! kind(any) does not make the lower-scored selector a strict superset. + ! CHECK-LABEL: func.func @_QMscoringPtest_any() + ! CHECK-NOT: fir.call @_QMscoringPlow + ! CHECK: fir.call @_QMscoringPhigh() + ! CHECK-NOT: fir.call @_QMscoringPlow + ! CHECK: return + subroutine test_any() + call base_any() + end subroutine + + subroutine base_sections() + !$omp declare variant(cpu) match(device={kind(cpu)}) + !$omp declare variant(vendor) & + !$omp& match(implementation={vendor(score(5): llvm)}) + end subroutine + + ! SECTION is a separator: the context depth is two, CPU scores 5, and the + ! vendor scores 6 whether the optional first SECTION is present or absent. + ! CHECK-LABEL: func.func @_QMscoringPtest_explicit_section() + ! CHECK: omp.parallel + ! CHECK: omp.sections + ! CHECK: omp.section + ! CHECK-NOT: fir.call @_QMscoringPcpu + ! CHECK: fir.call @_QMscoringPvendor() + ! CHECK-NOT: fir.call @_QMscoringPcpu + ! CHECK: return + subroutine test_explicit_section() + !$omp parallel sections + !$omp section + call base_sections() + !$omp end parallel sections + end subroutine + + ! CHECK-LABEL: func.func @_QMscoringPtest_implicit_section() + ! CHECK: omp.parallel + ! CHECK: omp.sections + ! CHECK: omp.section + ! CHECK-NOT: fir.call @_QMscoringPcpu + ! CHECK: fir.call @_QMscoringPvendor() + ! CHECK-NOT: fir.call @_QMscoringPcpu + ! CHECK: return + subroutine test_implicit_section() + !$omp parallel sections + call base_sections() + !$omp end parallel sections + end subroutine +end module + +! CHECK-LABEL: func.func @_QPtest_metadirective_any() +! CHECK-NOT: omp.taskyield +! CHECK: omp.barrier +! CHECK-NEXT: return +subroutine test_metadirective_any() + !$omp metadirective & + !$omp& when(implementation={vendor(score(100): llvm)}: barrier) & + !$omp& when(implementation={vendor(score(1): llvm)}, & + !$omp& device={kind(any)}: taskyield) +end subroutine diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h index 7849d32665994..18b7f884dfe37 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h @@ -21,6 +21,7 @@ #include "llvm/ADT/DenseMapInfo.h" #include "llvm/Frontend/OpenMP/OMPConstants.h" #include "llvm/Support/Compiler.h" +#include <optional> namespace llvm { class Triple; @@ -118,10 +119,37 @@ LLVM_ABI bool isValidTraitPropertyForTraitSetAndSelector(TraitProperty Property, TraitSelector Selector, TraitSet Set); -/// Variant match information describes the required traits and how they are -/// scored (via the ScoresMap). In addition, the required consturct nesting is -/// decribed as well. +/// Required traits, their scores, and the ordered construct selectors. +/// Unknown properties retain their identity and scores separately from the +/// enumerated properties so they can participate in matching extensions. struct VariantMatchInfo { + struct ISATrait { + TraitProperty Property; + StringRef Name; + + bool operator==(const ISATrait &Other) const { + return Property == Other.Property && Name == Other.Name; + } + }; + + struct UnknownTrait { + TraitSelector Selector; + StringRef Name; + std::optional<APInt> Score; + + // Subset checks compare property identity, not its score. + bool operator==(const UnknownTrait &Other) const { + return Selector == Other.Selector && Name == Other.Name; + } + }; + + /// Keep unknown properties inactive without losing their selector or score. + void addUnknownTrait(TraitSelector Selector, StringRef Name, + APInt *Score = nullptr) { + UnknownTraits.push_back( + {Selector, Name, Score ? std::optional<APInt>(*Score) : std::nullopt}); + } + /// Add the trait \p Property to the required trait set. \p RawString is the /// string we parsed and derived \p Property from. If \p Score is not null, it /// recorded as well. If \p Property is in the `construct` set it is recorded @@ -142,20 +170,27 @@ struct VariantMatchInfo { // Special handling for `device={isa(...)}` as we do not match the enum but // the raw string. - if (Property == TraitProperty::device_isa___ANY) - ISATraits.push_back(RawString); - if (Property == TraitProperty::target_device_isa___ANY) - ISATraits.push_back(RawString); + if (Property == TraitProperty::device_isa___ANY || + Property == TraitProperty::target_device_isa___ANY) + ISATraits.push_back({Property, RawString}); RequiredTraits.set(unsigned(Property)); if (Set == TraitSet::construct) ConstructTraits.push_back(Property); + if (getOpenMPContextTraitSelectorForProperty(Property) == + TraitSelector::user_condition) + UserCondition = RawString; } BitVector RequiredTraits = BitVector(unsigned(TraitProperty::Last) + 1); - SmallVector<StringRef, 8> ISATraits; + SmallVector<ISATrait, 8> ISATraits; SmallVector<TraitProperty, 8> ConstructTraits; SmallDenseMap<TraitProperty, APInt> ScoreMap; + SmallVector<UnknownTrait, 2> UnknownTraits; + /// Identity of a user condition expression, when the producer can retain + /// it. Applicability uses the folded property above; subset checks use this + /// identity to avoid conflating distinct dynamic expressions. + StringRef UserCondition; }; /// The context for a source location is made up of active property traits, @@ -175,12 +210,22 @@ struct OMPContext { ConstructTraits.push_back(Property); } + /// Record an enclosing construct that has no corresponding construct + /// selector property. Such constructs still contribute to the positions + /// and device-selector weights used during variant scoring. + void addUnknownConstruct() { + ConstructTraits.push_back(TraitProperty::invalid); + } + /// Hook for users to check if an ISA trait matches. The trait is described as /// the string that got parsed and it depends on the target and context if /// this matches or not. virtual bool matchesISATrait(StringRef) const { return false; } BitVector ActiveTraits = BitVector(unsigned(TraitProperty::Last) + 1); + /// Enclosing constructs in outermost-to-innermost order. An `invalid` + /// entry represents a construct that cannot itself appear in a construct + /// selector, but which still contributes to scoring depth. SmallVector<TraitProperty, 8> ConstructTraits; }; diff --git a/llvm/lib/Frontend/OpenMP/OMPContext.cpp b/llvm/lib/Frontend/OpenMP/OMPContext.cpp index f2cdd9bbaf5e4..798cef6598a53 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" @@ -138,13 +140,9 @@ OMPContext::OMPContext(bool IsDeviceCompilation, Triple TargetTriple, } } -/// Return true if \p C0 is a subset of \p C1. Note that both arrays are -/// expected to be sorted. -template <typename T> static bool isSubset(ArrayRef<T> C0, ArrayRef<T> C1) { -#ifdef EXPENSIVE_CHECKS - assert(llvm::is_sorted(C0) && llvm::is_sorted(C1) && - "Expected sorted arrays!"); -#endif +/// Return true if \p C0 is an ordered subsequence of \p C1. +template <typename T> +static bool isOrderedSubset(ArrayRef<T> C0, ArrayRef<T> C1) { if (C0.size() > C1.size()) return false; auto It0 = C0.begin(), End0 = C0.end(); @@ -157,24 +155,46 @@ template <typename T> static bool isSubset(ArrayRef<T> C0, ArrayRef<T> C1) { ++It1; continue; } - ++It0; + ++It1; } return true; } static bool isStrictSubset(const VariantMatchInfo &VMI0, const VariantMatchInfo &VMI1) { - // If all required traits are a strict subset and the ordered vectors storing - // the construct traits, we say it is a strict subset. Note that the latter - // relation is not required to be strict. - if (VMI0.RequiredTraits.count() >= VMI1.RequiredTraits.count()) + // kind(any) is equivalent to omitting the kind selector. + BitVector Traits0 = VMI0.RequiredTraits, Traits1 = VMI1.RequiredTraits; + for (TraitProperty Property : {TraitProperty::device_kind_any, + TraitProperty::target_device_kind_any}) { + Traits0.reset(unsigned(Property)); + Traits1.reset(unsigned(Property)); + } + size_t TraitCount0 = Traits0.count() + VMI0.UnknownTraits.size(); + size_t TraitCount1 = Traits1.count() + VMI1.UnknownTraits.size(); + if (TraitCount0 > TraitCount1) return false; - for (unsigned Bit : VMI0.RequiredTraits.set_bits()) - if (!VMI1.RequiredTraits.test(Bit)) + for (unsigned Bit : Traits0.set_bits()) + if (!Traits1.test(Bit)) + return false; + for (const auto &Trait : VMI0.UnknownTraits) + if (!llvm::is_contained(VMI1.UnknownTraits, Trait)) return false; - if (!isSubset<TraitProperty>(VMI0.ConstructTraits, VMI1.ConstructTraits)) + for (const auto &Trait : VMI0.ISATraits) + if (!llvm::is_contained(VMI1.ISATraits, Trait)) + return false; + bool HasAdditionalISATrait = + llvm::any_of(VMI1.ISATraits, [&](const auto &Trait) { + return !llvm::is_contained(VMI0.ISATraits, Trait); + }); + if (!VMI0.UserCondition.empty() && VMI0.UserCondition != VMI1.UserCondition) return false; - return true; + if (!isOrderedSubset<TraitProperty>(VMI0.ConstructTraits, + VMI1.ConstructTraits)) + return false; + // RequiredTraits is a bit vector, so repeated construct properties only + // make the ordered vector strict. + return TraitCount0 < TraitCount1 || HasAdditionalISATrait || + VMI0.ConstructTraits.size() < VMI1.ConstructTraits.size(); } static int @@ -197,23 +217,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({ @@ -229,6 +247,9 @@ isVariantApplicableInContextHelper(const VariantMatchInfo &VMI, return false; }; + if (!VMI.UnknownTraits.empty() && !HandleTrait(TraitProperty::invalid, false)) + return false; + for (unsigned Bit : VMI.RequiredTraits.set_bits()) { TraitProperty Property = TraitProperty(Bit); if (DeviceOrImplementationSetOnly && @@ -247,39 +268,38 @@ isVariantApplicableInContextHelper(const VariantMatchInfo &VMI, // We overwrite the isa trait as it is actually up to the OMPContext hook to // check the raw string(s). - if (Property == TraitProperty::device_isa___ANY) - IsActiveTrait = llvm::all_of(VMI.ISATraits, [&](StringRef RawString) { - return Ctx.matchesISATrait(RawString); - }); - if (Property == TraitProperty::target_device_isa___ANY) - IsActiveTrait = llvm::all_of(VMI.ISATraits, [&](StringRef RawString) { - return Ctx.matchesISATrait(RawString); + if (Property == TraitProperty::device_isa___ANY || + Property == TraitProperty::target_device_isa___ANY) + IsActiveTrait = llvm::all_of(VMI.ISATraits, [&](const auto &Trait) { + return Trait.Property != Property || Ctx.matchesISATrait(Trait.Name); }); - if (std::optional<bool> Result = HandleTrait(Property, IsActiveTrait)) - return *Result; + if (!HandleTrait(Property, IsActiveTrait)) + return false; } if (!DeviceOrImplementationSetOnly) { - // We could use isSubset 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"); @@ -289,11 +309,29 @@ isVariantApplicableInContextHelper(const VariantMatchInfo &VMI, // TODO: Verify SIMD } - assert(isSubset<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"); @@ -313,118 +351,116 @@ 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. Unknown + // properties retain these scores even though they never match. + 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); } + for (const auto &Trait : VMI.UnknownTraits) + AddSelectorScore(Trait.Selector, Trait.Score ? &*Trait.Score : nullptr); - 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); - int BestVMIIdx = -1; - const VariantMatchInfo *BestVMI = nullptr; - + SmallVector<std::optional<APInt>, 4> Scores(VMIs.size()); for (unsigned u = 0, e = VMIs.size(); u < e; ++u) { const VariantMatchInfo &VMI = VMIs[u]; SmallVector<unsigned, 8> ConstructMatches; - // If the variant is not applicable its not the best. + // Inapplicable variants do not participate in scoring or subset checks. if (!isVariantApplicableInContextHelper( VMI, Ctx, &ConstructMatches, /* DeviceOrImplementationSetOnly */ false)) continue; - // Check if its clearly not the best. - APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches); - if (Score.ult(BestScore)) + Scores[u] = getVariantMatchScore(VMI, Ctx, ConstructMatches); + } + + // A compatible selector that is a strict subset of another compatible + // selector has score zero, irrespective of their scores before this step. + // Apply this rule globally before choosing the maximum. + for (unsigned u = 0, e = VMIs.size(); u < e; ++u) { + if (!Scores[u]) continue; - // Equal score need subset checks. - if (Score.eq(BestScore)) { - // Strict subset are never best. - if (isStrictSubset(VMI, *BestVMI)) - continue; - // Same score and the current best is no strict subset so we keep it. - if (!isStrictSubset(*BestVMI, VMI)) - continue; + for (unsigned v = 0; v < e; ++v) { + if (u != v && Scores[v] && isStrictSubset(VMIs[u], VMIs[v])) { + Scores[u] = APInt(1, 0); + break; + } } - // New best found. - BestVMI = &VMI; + } + + APInt BestScore(1, 0); + int BestVMIIdx = -1; + for (unsigned u = 0, e = VMIs.size(); u < e; ++u) { + if (!Scores[u]) + continue; + const APInt &Score = *Scores[u]; + unsigned Width = std::max(Score.getBitWidth(), BestScore.getBitWidth()); + if (BestVMIIdx >= 0 && + !BestScore.zextOrTrunc(Width).ult(Score.zextOrTrunc(Width))) + continue; BestVMIIdx = u; BestScore = Score; } diff --git a/llvm/unittests/Frontend/OpenMPContextTest.cpp b/llvm/unittests/Frontend/OpenMPContextTest.cpp index f9683ae56e933..ae81615f3e0d6 100644 --- a/llvm/unittests/Frontend/OpenMPContextTest.cpp +++ b/llvm/unittests/Frontend/OpenMPContextTest.cpp @@ -314,7 +314,464 @@ 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); + + VariantMatchInfo Parallel; + Parallel.addTrait(TraitProperty::construct_parallel_parallel, ""); + + VariantMatchInfo MatchAny; + MatchAny.addTrait(TraitProperty::construct_target_target, ""); + MatchAny.addTrait(TraitProperty::construct_parallel_parallel, ""); + MatchAny.addTrait(TraitProperty::implementation_extension_match_any, ""); + + // The unadjusted scores are 1 + 1 + 2 for MATCH_ANY and 1 + 2 for PARALLEL. + // PARALLEL is also a strict subset, so its final score is zero. + SmallVector<VariantMatchInfo, 2> Candidates{Parallel, 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 MATCH_ANY applicable, but the absent construct + // must not add to its score. The scored candidate wins by 2 to 1. + SmallVector<VariantMatchInfo, 2> Candidates{MatchAny, Scored}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, NoConstruct), 1); +} + +TEST_F(OpenMPContextTest, ScoringUnknownProperty) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + VariantMatchInfo Unknown; + Unknown.addUnknownTrait(TraitSelector::implementation_vendor, "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}; + // An unscored vendor contributes zero, so lexical 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, ScoringUnknownDeviceProperties) { + const TraitSelector Selectors[] = { + TraitSelector::device_kind, TraitSelector::device_arch, + TraitSelector::target_device_kind, TraitSelector::target_device_arch}; + for (unsigned Depth : {0u, 2u, 64u}) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + for (unsigned I = 0; I < Depth; ++I) + Context.addTrait(TraitProperty::construct_parallel_parallel); + for (unsigned I = 0; I < 4; ++I) { + SCOPED_TRACE(Depth); + SCOPED_TRACE(I); + VariantMatchInfo Unknown; + Unknown.addUnknownTrait(Selectors[I], "unknown"); + EXPECT_FALSE(isVariantApplicableInContext(Unknown, Context)); + Unknown.addTrait(TraitProperty::implementation_extension_match_any, ""); + EXPECT_FALSE(isVariantApplicableInContext(Unknown, Context)); + Unknown.addTrait(TraitProperty::user_condition_true, ""); + EXPECT_TRUE(isVariantApplicableInContext(Unknown, Context)); + + APInt Weight = APInt::getOneBitSet(128, Depth + I % 2); + APInt Score = Weight - 1; + VariantMatchInfo Scored; + Scored.addTrait(TraitProperty::implementation_vendor_llvm, "", &Score); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, Unknown}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + Candidates[0].addTrait(TraitProperty::implementation_vendor_llvm, "", + &Weight); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + + // match_none also retains the weight of an inactive selector. + Candidates[1] = VariantMatchInfo(); + Candidates[1].addUnknownTrait(Selectors[I], "unknown"); + Candidates[1].addTrait(TraitProperty::implementation_extension_match_none, + ""); + EXPECT_TRUE(isVariantApplicableInContext(Candidates[1], Context)); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + Candidates[0].addTrait(TraitProperty::implementation_vendor_llvm, "", + &Score); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + } + } +} + +TEST_F(OpenMPContextTest, ScoringUnknownSelectors) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + APInt VendorScore(64, 10), ExtensionScore(64, 1), ConditionScore(64, 5); + for (bool Reverse : {false, true}) { + SCOPED_TRACE(Reverse); + VariantMatchInfo Unknown; + Unknown.addTrait(TraitProperty::implementation_extension_match_any, "", + &ExtensionScore); + Unknown.addTrait(TraitProperty::user_condition_true, "", &ConditionScore); + Unknown.addUnknownTrait(TraitSelector::implementation_vendor, "vendor", + &VendorScore); + Unknown.addUnknownTrait(TraitSelector::implementation_extension, + "extension", &ExtensionScore); + if (Reverse) + std::swap(Unknown.UnknownTraits[0], Unknown.UnknownTraits[1]); + + VariantMatchInfo Scored; + APInt Score(64, 15); + Scored.addTrait(TraitProperty::implementation_vendor_llvm, "", &Score); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, Unknown}; + // Each selector contributes once: 1 + 10 + 1 + 5 = 17. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + Score = APInt(64, 16); + Candidates[0].addTrait(TraitProperty::implementation_vendor_llvm, "", + &Score); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + } +} + +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); + + Candidates[1].addTrait(TraitProperty::implementation_extension_match_any, ""); + Candidates[1].addUnknownTrait(TraitSelector::device_kind, "unknown"); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); +} + +TEST_F(OpenMPContextTest, UnknownPropertySubsets) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + VariantMatchInfo First, Second; + First.addTrait(TraitProperty::implementation_extension_match_none, ""); + First.addUnknownTrait(TraitSelector::implementation_vendor, "first"); + Second.addTrait(TraitProperty::implementation_extension_match_none, ""); + Second.addUnknownTrait(TraitSelector::implementation_vendor, "second"); + Second.addUnknownTrait(TraitSelector::implementation_extension, "extension"); + SmallVector<VariantMatchInfo, 2> Candidates{First, Second}; + // Distinct unknown properties do not form a subset just because their + // selectors are the same. The scores tie, so the first candidate wins. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + Candidates[1].UnknownTraits[0].Name = "first"; + // Now the first candidate is a strict subset of the second. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); +} + +TEST_F(OpenMPContextTest, StrictSubsetScoreIsZero) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + APInt HighScore(64, 100), LowScore(64, 1); + VariantMatchInfo Subset, Superset; + Subset.addTrait(TraitProperty::implementation_vendor_llvm, "", &HighScore); + Superset.addTrait(TraitProperty::implementation_vendor_llvm, "", &LowScore); + Superset.addTrait(TraitProperty::user_condition_true, ""); + + // The strict subset's score is zero even though its raw score is higher. + SmallVector<VariantMatchInfo, 2> Candidates{Subset, Superset}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + std::swap(Candidates[0], Candidates[1]); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); +} + +TEST_F(OpenMPContextTest, KindAnyDoesNotAffectSubsets) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + APInt HighScore(64, 100), LowScore(64, 1); + VariantMatchInfo High, Low; + High.addTrait(TraitProperty::implementation_vendor_llvm, "", &HighScore); + Low.addTrait(TraitProperty::implementation_vendor_llvm, "", &LowScore); + Low.addTrait(TraitProperty::device_kind_any, ""); + + SmallVector<VariantMatchInfo, 2> Candidates{High, Low}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + std::swap(Candidates[0], Candidates[1]); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + + // Ignoring any must also work when it appears in the potential subset. + High.addTrait(TraitProperty::device_kind_any, ""); + Low.RequiredTraits.reset(unsigned(TraitProperty::device_kind_any)); + Low.addTrait(TraitProperty::device_kind_host, ""); + Candidates = {High, Low}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); +} + +TEST_F(OpenMPContextTest, DistinctUserConditionsAreNotSubsets) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + APInt HighScore(64, 100); + VariantMatchInfo High, MoreTraits; + High.addTrait(TraitProperty::user_condition_true, "high", &HighScore); + MoreTraits.addTrait(TraitProperty::user_condition_true, "low"); + MoreTraits.addTrait(TraitProperty::device_kind_host, ""); + + // The condition expressions are different properties. HIGH is not a strict + // subset of MORE_TRAITS and retains its explicit score. + SmallVector<VariantMatchInfo, 2> Candidates{MoreTraits, High}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); +} + +TEST_F(OpenMPContextTest, RepeatedConstructStrictSubsetScoreIsZero) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + Context.addTrait(TraitProperty::construct_parallel_parallel); + Context.addTrait(TraitProperty::construct_parallel_parallel); + APInt HighScore(64, 100), LowScore(64, 1); + VariantMatchInfo Subset, Superset; + Subset.addTrait(TraitProperty::implementation_vendor_llvm, "", &HighScore); + Subset.addTrait(TraitProperty::construct_parallel_parallel, ""); + Superset.addTrait(TraitProperty::implementation_vendor_llvm, "", &LowScore); + Superset.addTrait(TraitProperty::construct_parallel_parallel, ""); + Superset.addTrait(TraitProperty::construct_parallel_parallel, ""); + + // Repeated constructs are retained only in the ordered vector, not the bit + // vector. They must still make the second selector a strict superset. + SmallVector<VariantMatchInfo, 2> Candidates{Subset, Superset}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); +} + +TEST_F(OpenMPContextTest, DifferentConstructOrderIsNotSubset) { + 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); + + APInt HighScore(64, 100), LowScore(64, 1); + VariantMatchInfo High, Low; + High.addTrait(TraitProperty::implementation_vendor_llvm, "", &HighScore); + High.addTrait(TraitProperty::construct_parallel_parallel, ""); + High.addTrait(TraitProperty::construct_for_for, ""); + Low.addTrait(TraitProperty::implementation_vendor_llvm, "", &LowScore); + Low.addTrait(TraitProperty::user_condition_true, ""); + Low.addTrait(TraitProperty::construct_for_for, ""); + Low.addTrait(TraitProperty::construct_parallel_parallel, ""); + Low.addTrait(TraitProperty::construct_simd_simd, ""); + + // HIGH's construct sequence is not a subsequence of LOW's, despite each + // distinct construct property being present in LOW. + SmallVector<VariantMatchInfo, 2> Candidates{High, Low}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + std::swap(Candidates[0], Candidates[1]); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); +} + +TEST_F(OpenMPContextTest, DifferentISAPropertiesAreNotSubsets) { + struct ISAContext : OMPContext { + using OMPContext::OMPContext; + bool matchesISATrait(StringRef) const override { return true; } + }; + ISAContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + APInt HighScore(64, 100), LowScore(64, 1); + VariantMatchInfo High, Low; + High.addTrait(TraitProperty::implementation_vendor_llvm, "", &HighScore); + High.addTrait(TraitProperty::device_isa___ANY, "sse2"); + Low.addTrait(TraitProperty::implementation_vendor_llvm, "", &LowScore); + Low.addTrait(TraitProperty::device_isa___ANY, "avx"); + Low.addTrait(TraitProperty::user_condition_true, ""); + + SmallVector<VariantMatchInfo, 2> Candidates{High, Low}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0); + std::swap(Candidates[0], Candidates[1]); + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + + // Once the ISA properties are identical, HIGH is a strict subset of LOW. + Candidates[0].ISATraits[0].Name = "sse2"; + 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, ScoringUnknownConstructs) { + OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1); + Context.addTrait(TraitProperty::construct_parallel_parallel); + Context.addUnknownConstruct(); + + APInt Score(64, 3); + VariantMatchInfo Scored, Device; + Scored.addTrait(TraitProperty::implementation_vendor_llvm, "", &Score); + Device.addTrait(TraitProperty::device_kind_cpu, ""); + SmallVector<VariantMatchInfo, 2> Candidates{Scored, Device}; + // Both enclosing constructs contribute to the device weight: 1 + 2^2. + EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1); + + OMPContext UnknownThenParallel(false, Triple("x86_64-unknown-linux"), + Triple(), -1); + UnknownThenParallel.addUnknownConstruct(); + UnknownThenParallel.addTrait(TraitProperty::construct_parallel_parallel); + APInt LowerScore(64, 1); + VariantMatchInfo LowerScored, Parallel; + LowerScored.addTrait(TraitProperty::implementation_vendor_llvm, "", + &LowerScore); + Parallel.addTrait(TraitProperty::construct_parallel_parallel, ""); + // The unknown construct occupies position one, so PARALLEL scores 1 + 2. + EXPECT_TRUE(isVariantApplicableInContext(Parallel, UnknownThenParallel)); + Candidates = {LowerScored, Parallel}; + EXPECT_EQ(getBestVariantMatchForContext(Candidates, UnknownThenParallel), 1); +} + +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); + } } } // namespace _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
