[flang-commits] [clang] [flang] [llvm] [clang][flang][OpenMP] Fix context selector matching and scoring (PR #224431)
via flang-commits
flang-commits at lists.llvm.org
Wed Sep 23 22:39:29 PDT 2026
https://github.com/chichunchen updated https://github.com/llvm/llvm-project/pull/224431
>From 433b44353c08e386c62c431b38a734aff04e64c3 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Wed, 16 Sep 2026 16:48:23 -0500
Subject: [PATCH 1/2] [clang][flang][OpenMP] Fix context selector matching and
scoring
Incorrect construct contexts and selector scoring can select the wrong
DECLARE VARIANT function or METADIRECTIVE replacement.
Build construct contexts in source order rather than emitted MLIR order:
Source: teams -> distribute -> parallel -> do
MLIR: teams -> parallel -> distribute -> do
Count executable constructs without selectable traits, omit informational
directives, and start the context at the innermost TARGET.
Compute device weights from the enclosing context depth, choose the
highest-scoring complete ordered construct match, and count each selector's
score once using arbitrary-width arithmetic.
Inside a single PARALLEL region, device={kind(cpu)} previously scored 2,
tying construct={parallel}. Using the enclosing context depth gives the
device selector a score of 3, so it correctly wins.
Zero strict-subset scores before selecting the best applicable candidate.
For example, when both selectors match:
A: implementation={vendor(score(100): llvm)}
B: implementation={vendor(score(1): llvm)}, user={condition(1)}
A initially scores 101 and B scores 2. Since A is a strict subset of B,
A's final score is zero and B wins. Previously, subset checks applied
only to equal scores.
Preserve property and condition identity in subset comparisons, including
across Clang AST serialization. Keep Flang runtime and folded conditions
consistent: false conditions retain their identity, contribute no score
under match_any, and retain their score under match_none. Apply the same
ranking rules to implicit NOTHING replacements.
Add shared matcher unit tests and Clang/Flang regressions, including PCH
condition identity and static/runtime ranking.
This provides the shared matching and scoring prerequisite for #219014.
Assisted with Codex.
---
clang/include/clang/AST/OpenMPClause.h | 2 +-
.../clang/Serialization/ASTRecordWriter.h | 2 +
clang/lib/AST/OpenMPClause.cpp | 30 +-
clang/lib/Parse/ParseOpenMP.cpp | 6 +-
clang/lib/Sema/SemaOpenMP.cpp | 50 +-
clang/lib/Serialization/ASTReader.cpp | 4 +-
clang/lib/Serialization/ASTWriter.cpp | 4 +-
.../declare_variant_condition_identity_pch.c | 57 ++
.../declare_variant_device_kind_codegen.cpp | 45 +-
clang/test/OpenMP/declare_variant_scoring.c | 289 ++++++++++
clang/test/OpenMP/dispatch_variant_matching.c | 43 ++
flang/include/flang/Lower/AbstractConverter.h | 3 +
flang/include/flang/Semantics/openmp-utils.h | 9 +
flang/lib/Lower/Bridge.cpp | 6 +
flang/lib/Lower/OpenMP/Atomic.cpp | 3 +
flang/lib/Lower/OpenMP/OpenMP.cpp | 16 +-
flang/lib/Lower/OpenMP/Utils.cpp | 67 ++-
flang/lib/Lower/OpenMP/Utils.h | 23 +-
flang/lib/Semantics/openmp-utils.cpp | 103 +++-
.../OpenMP/declare-variant-construct.f90 | 169 +++++-
.../OpenMP/declare-variant-source-context.f90 | 113 ++++
.../OpenMP/metadirective-device-kind.f90 | 17 +
.../OpenMP/metadirective-implementation.f90 | 12 +
.../OpenMP/metadirective-source-context.f90 | 68 +++
.../OpenMP/metadirective-target-boundary.f90 | 62 ++
.../OpenMP/metadirective-user-static.f90 | 81 +++
.../test/Lower/OpenMP/metadirective-user.f90 | 306 +++++++---
.../OpenMP/variant-scoring-any-sections.f90 | 81 +++
.../include/llvm/Frontend/OpenMP/OMPContext.h | 60 +-
llvm/lib/Frontend/OpenMP/OMPContext.cpp | 304 +++++-----
llvm/unittests/Frontend/OpenMPContextTest.cpp | 531 +++++++++++++++++-
31 files changed, 2239 insertions(+), 327 deletions(-)
create mode 100644 clang/test/OpenMP/declare_variant_condition_identity_pch.c
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/declare-variant-source-context.f90
create mode 100644 flang/test/Lower/OpenMP/metadirective-source-context.f90
create mode 100644 flang/test/Lower/OpenMP/metadirective-target-boundary.f90
create mode 100644 flang/test/Lower/OpenMP/metadirective-user-static.f90
create mode 100644 flang/test/Lower/OpenMP/variant-scoring-any-sections.f90
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index ec84f10956ff43..18f1b404445b8c 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -10057,7 +10057,7 @@ struct OMPTraitProperty {
llvm::omp::TraitProperty Kind = llvm::omp::TraitProperty::invalid;
/// The raw string as we parsed it. This is needed for the `isa` trait set
- /// (which accepts anything) and (later) extensions.
+ /// (which accepts anything), user-condition identity, and extensions.
StringRef RawString;
};
diff --git a/clang/include/clang/Serialization/ASTRecordWriter.h b/clang/include/clang/Serialization/ASTRecordWriter.h
index 9849ea6b395ab6..6e52b7473c3692 100644
--- a/clang/include/clang/Serialization/ASTRecordWriter.h
+++ b/clang/include/clang/Serialization/ASTRecordWriter.h
@@ -189,6 +189,8 @@ class ASTRecordWriter
Record->push_back(Value);
}
+ void writeString(StringRef Value) { Writer->AddString(Value, *Record); }
+
void writeUnsignedOrNone(UnsignedOrNone Value) {
Record->push_back(Value.toInternalRepresentation());
}
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 2061d5395ac658..077b81c47cae49 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -3130,9 +3130,10 @@ void OMPTraitInfo::getAsVariantMatchInfo(ASTContext &ASTCtx,
Selector.ScoreOrCondition->getIntegerConstantExpr(ASTCtx))
VMI.addTrait(CondVal->isZero() ? TraitProperty::user_condition_false
: TraitProperty::user_condition_true,
- "<condition>");
+ Selector.Properties.front().RawString);
else
- VMI.addTrait(TraitProperty::user_condition_false, "<condition>");
+ VMI.addTrait(TraitProperty::user_condition_false,
+ Selector.Properties.front().RawString);
continue;
}
@@ -3302,8 +3303,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/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 0ce484c5e907d3..6f558df8e8b851 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -1150,8 +1150,12 @@ void Parser::parseOMPContextSelector(
if (!Condition.isUsable())
return FinishSelector();
TISelector.ScoreOrCondition = Condition.get();
+ StringRef ConditionText = Lexer::getSourceText(
+ CharSourceRange::getTokenRange(Condition.get()->getSourceRange()),
+ PP.getSourceManager(), PP.getLangOpts());
TISelector.Properties.push_back(
- {TraitProperty::user_condition_unknown, "<condition>"});
+ {TraitProperty::user_condition_unknown,
+ ConditionText.empty() ? "<condition>" : ConditionText});
return;
}
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 2e4d9f2f82f0b7..a7eb0079a28c16 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -4448,16 +4448,46 @@ 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)) {
+ if (llvm::omp::getDirectiveCategory(Leaf) ==
+ llvm::omp::Category::Informational)
+ continue;
+ 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/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index a9c230d767c509..8910b28e6c1659 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -13135,8 +13135,10 @@ OMPTraitInfo *ASTRecordReader::readOMPTraitInfo() {
if (readBool())
Selector.ScoreOrCondition = readExprRef();
Selector.Properties.resize(readUInt32());
- for (auto &Property : Selector.Properties)
+ for (auto &Property : Selector.Properties) {
Property.Kind = readEnum<llvm::omp::TraitProperty>();
+ Property.RawString = getContext().backupStr(readString());
+ }
}
}
return &TI;
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 944e68ea6481d9..a2db5b1a9bd1e6 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8978,8 +8978,10 @@ void ASTRecordWriter::writeOMPTraitInfo(const OMPTraitInfo *TI) {
if (Selector.ScoreOrCondition)
writeExprRef(Selector.ScoreOrCondition);
writeUInt32(Selector.Properties.size());
- for (const auto &Property : Selector.Properties)
+ for (const auto &Property : Selector.Properties) {
writeEnum(Property.Kind);
+ writeString(Property.RawString);
+ }
}
}
}
diff --git a/clang/test/OpenMP/declare_variant_condition_identity_pch.c b/clang/test/OpenMP/declare_variant_condition_identity_pch.c
new file mode 100644
index 00000000000000..3b0acd9f8f1f6b
--- /dev/null
+++ b/clang/test/OpenMP/declare_variant_condition_identity_pch.c
@@ -0,0 +1,57 @@
+// RUN: split-file %s %t
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -x c-header %t/conditions.h \
+// RUN: -emit-pch -o %t/conditions-c.pch
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -x c %t/use.c \
+// RUN: -include-pch %t/conditions-c.pch -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -x c++-header %t/conditions.h \
+// RUN: -emit-pch -o %t/conditions-cxx.pch
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -x c++ %t/use.c \
+// RUN: -include-pch %t/conditions-cxx.pch -emit-llvm -o - | FileCheck %s
+
+// Verify that serialization preserves the source identity of folded user
+// conditions. Distinct conditions do not form a subset relationship, while
+// identical conditions do.
+
+// CHECK-LABEL: define{{.*}} void @test_conditions()
+// CHECK: call void @condition_high_variant()
+// CHECK-NEXT: call void @condition_low_variant()
+// CHECK: ret void
+
+//--- conditions.h
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+void condition_high_variant(void);
+void condition_low_variant(void);
+
+#pragma omp declare variant(condition_high_variant) \
+ match(implementation = {vendor(score(100) : llvm)}, user = {condition(1)})
+#pragma omp declare variant(condition_low_variant) \
+ match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
+ user = {condition(2)})
+void distinct_condition_base(void);
+
+#pragma omp declare variant(condition_high_variant) \
+ match(implementation = {vendor(score(100) : llvm)}, user = {condition(1)})
+#pragma omp declare variant(condition_low_variant) \
+ match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
+ user = {condition(1)})
+void identical_condition_base(void);
+
+#ifdef __cplusplus
+}
+#endif
+
+//--- use.c
+#ifdef __cplusplus
+extern "C"
+#endif
+void test_conditions(void) {
+ distinct_condition_base();
+ identical_condition_base();
+}
diff --git a/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp b/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp
index 9335df10f957c9..d66d265caa8ebb 100644
--- a/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp
+++ b/clang/test/OpenMP/declare_variant_device_kind_codegen.cpp
@@ -82,21 +82,29 @@
// Verify no unexpected global symbol aliasing
// CHECK-NOT: @{{[^ ]+}} = {{.*}}alias
-
-// CHECK-NOT: ret i32 {{1|4|81|84}}
-// CHECK-DAG: declare {{.*}}i32 @_Z5bazzzv()
-// CHECK-DAG: define {{.*}}i32 @_Z3bazv()
-// CHECK-DAG: ret i32 2
-// CHECK-DAG: ret i32 3
-// CHECK-DAG: ret i32 5
-// CHECK-DAG: ret i32 6
-// CHECK-DAG: ret i32 7
-// CHECK-DAG: ret i32 82
-// CHECK-DAG: ret i32 83
-// CHECK-DAG: ret i32 85
-// CHECK-DAG: ret i32 86
-// CHECK-DAG: ret i32 87
-// CHECK-NOT: ret i32 {{4|81|84}}
+// CHECK-LABEL: define {{.*}}i32 @main()
+// CHECK: call {{.*}}i32 @_Z3foov()
+// CHECK: call {{.*}}i32 @_Z5bazzzv()
+// CHECK: call {{.*}}i32 @_Z4testv()
+// CHECK: call {{.*}}i32 @_ZL10stat_used_v()
+// CHECK: ret i32
+
+// CHECK-LABEL: define {{.*}}void @_Z3xxxv()
+// CHECK: call {{.*}}i32 @_ZN12SpecialFuncs7method_Ev(
+// CHECK: call {{.*}}i32 @_ZN16SpecSpecialFuncs7method_Ev(
+// CHECK: ret void
+
+// CHECK-LABEL: define {{.*}}i32 @_Z6int_fnv()
+// CHECK-DAG: call {{.*}}i32 @_Z5prio1v()
+// CHECK-DAG: call {{.*}}i32 @_ZL5prio2v()
+// CHECK: ret i32
+
+// CHECK-LABEL: define {{.*}}i32 @_Z18fn_linkage_variantv()
+// CHECK: ret i32 85
+// CHECK-LABEL: define {{.*}}i32 @fn_linkage_variant1()
+// CHECK: ret i32 86
+// CHECK-LABEL: define {{.*}}i32 @_Z3fn2v()
+// CHECK: ret i32 87
#ifndef HEADER
#define HEADER
@@ -108,8 +116,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
@@ -199,7 +208,7 @@ static int prio4() { return 84; }
#pragma omp declare variant(prio3) match(device = {kind(SUBSET)})
static int prio1_() { return 1; }
-int int_fn() { return prio1_(); }
+int int_fn() { return prio_() + prio1_(); }
int fn_linkage_variant() { return 85; }
extern "C" {
diff --git a/clang/test/OpenMP/declare_variant_scoring.c b/clang/test/OpenMP/declare_variant_scoring.c
new file mode 100644
index 00000000000000..c924dbaf346246
--- /dev/null
+++ b/clang/test/OpenMP/declare_variant_scoring.c
@@ -0,0 +1,289 @@
+// 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 arch_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);
+
+#pragma omp declare variant(arch_variant) match(device = {arch(x86_64)})
+#pragma omp declare variant(scored_variant) \
+ match(implementation = {vendor(score(3) : llvm)})
+void source_depth_base(void);
+
+int arch_value(void);
+int scored_value(void);
+#pragma omp declare variant(arch_value) match(device = {arch(x86_64)})
+#pragma omp declare variant(scored_value) \
+ match(implementation = {vendor(score(3) : llvm)})
+int source_depth_value(void);
+
+void condition_high_variant(void);
+void condition_low_variant(void);
+#pragma omp declare variant(condition_high_variant) \
+ match(implementation = {vendor(score(100) : llvm)}, user = {condition(1)})
+#pragma omp declare variant(condition_low_variant) \
+ match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
+ user = {condition(2)})
+void distinct_condition_base(void);
+
+#pragma omp declare variant(condition_high_variant) \
+ match(implementation = {vendor(score(100) : llvm)}, user = {condition(1)})
+#pragma omp declare variant(condition_low_variant) \
+ match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
+ user = {condition(1)})
+void identical_condition_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()
+
+// Executable loop transformations contribute to the construct depth. At
+// depth one, ARCH scores 5 and beats the vendor variant's score of 4.
+void tile_depth(void) {
+#pragma omp tile sizes(2)
+ for (int i = 0; i < 4; ++i)
+ source_depth_base();
+}
+// CHECK-LABEL: define{{.*}} void @tile_depth
+// CHECK: call void @arch_variant()
+// CHECK: ret void
+
+void unroll_depth(void) {
+#pragma omp unroll partial(2)
+ for (int i = 0; i < 4; ++i)
+ source_depth_base();
+}
+// CHECK-LABEL: define{{.*}} void @unroll_depth
+// CHECK: call void @arch_variant()
+// CHECK: ret void
+
+// ATOMIC is an executable construct and contributes one context position.
+void atomic_depth(int *x) {
+#pragma omp atomic update
+ *x += source_depth_value();
+}
+// CHECK-LABEL: define{{.*}} void @atomic_depth
+// CHECK: call i32 @arch_value()
+// CHECK: ret void
+
+// ASSUME is informational and does not add a position. Inside PARALLEL, CPU
+// scores 3 and loses to the vendor variant's score of 4.
+void assume_context(void) {
+#pragma omp parallel
+ {
+#pragma omp assume no_openmp_routines
+ { task_depth_base(); }
+ }
+}
+// CHECK-LABEL: define internal void @assume_context.omp_outlined
+// CHECK: call void @scored_variant()
+// CHECK: ret void
+
+// Distinct folded user conditions do not create a subset relationship.
+void distinct_conditions(void) { distinct_condition_base(); }
+// CHECK-LABEL: define{{.*}} void @distinct_conditions
+// CHECK: call void @condition_high_variant()
+// CHECK: ret void
+
+// Identical conditions allow the first selector to be a strict subset.
+void identical_conditions(void) { identical_condition_base(); }
+// CHECK-LABEL: define{{.*}} void @identical_conditions
+// CHECK: call void @condition_low_variant()
+// CHECK: ret void
+
+// 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 00000000000000..e61f719e77293d
--- /dev/null
+++ b/clang/test/OpenMP/dispatch_variant_matching.c
@@ -0,0 +1,43 @@
+// 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());
+}
diff --git a/flang/include/flang/Lower/AbstractConverter.h b/flang/include/flang/Lower/AbstractConverter.h
index ae246d3188bd88..eb17d03d458535 100644
--- a/flang/include/flang/Lower/AbstractConverter.h
+++ b/flang/include/flang/Lower/AbstractConverter.h
@@ -337,6 +337,9 @@ class AbstractConverter {
/// Get the converter's current scope
virtual const Fortran::semantics::Scope &getCurrentScope() = 0;
+ /// Get the PFT evaluation currently being lowered.
+ virtual pft::Evaluation &getCurrentEvaluation() = 0;
+
//===--------------------------------------------------------------------===//
// FIR/MLIR
//===--------------------------------------------------------------------===//
diff --git a/flang/include/flang/Semantics/openmp-utils.h b/flang/include/flang/Semantics/openmp-utils.h
index 769140a3147edf..bdab224907346f 100644
--- a/flang/include/flang/Semantics/openmp-utils.h
+++ b/flang/include/flang/Semantics/openmp-utils.h
@@ -242,6 +242,15 @@ class OmpVariantMatchContext : public llvm::omp::OMPContext {
std::string features_;
};
+/// Append the construct context contributed by \p directive.
+///
+/// Combined and composite directives contribute their leaf constructs in
+/// source order. Informational directives do not contribute. Entering a TARGET
+/// construct discards the outer context, as required by the OpenMP
+/// construct-set definition.
+void AppendDirectiveContextTraits(llvm::omp::Directive directive,
+ llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits);
+
struct MetadirectiveCandidate {
MetadirectiveCandidate(const parser::OmpDirectiveSpecification *spec,
llvm::omp::VariantMatchInfo vmi, bool isExplicit,
diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp
index 784a87cc619979..f1c127e01513ae 100644
--- a/flang/lib/Lower/Bridge.cpp
+++ b/flang/lib/Lower/Bridge.cpp
@@ -1122,6 +1122,8 @@ class FirConverter : public Fortran::lower::AbstractConverter {
void genEval(Fortran::lower::pft::Evaluation &eval,
bool unstructuredContext) override final {
+ Fortran::lower::pft::Evaluation *previousEval{evalPtr};
+ llvm::scope_exit restoreEval([&]() { evalPtr = previousEval; });
genFIR(eval, unstructuredContext);
}
@@ -1227,6 +1229,10 @@ class FirConverter : public Fortran::lower::AbstractConverter {
return bridge.getSemanticsContext().FindScope(currentPosition);
}
+ Fortran::lower::pft::Evaluation &getCurrentEvaluation() override final {
+ return getEval();
+ }
+
fir::FirOpBuilder &getFirOpBuilder() override final {
CHECK(builder && "builder is not set before calling getFirOpBuilder");
return *builder;
diff --git a/flang/lib/Lower/OpenMP/Atomic.cpp b/flang/lib/Lower/OpenMP/Atomic.cpp
index 5a38c97ee77a2f..4a73fae32b9346 100644
--- a/flang/lib/Lower/OpenMP/Atomic.cpp
+++ b/flang/lib/Lower/OpenMP/Atomic.cpp
@@ -7,6 +7,7 @@
//===----------------------------------------------------------------------===//
#include "Atomic.h"
+#include "Utils.h"
#include "flang/Evaluate/expression.h"
#include "flang/Evaluate/fold.h"
#include "flang/Evaluate/tools.h"
@@ -559,6 +560,8 @@ void Fortran::lower::omp::lowerAtomic(
fir::FirOpBuilder &builder = converter.getFirOpBuilder();
const parser::OmpDirectiveSpecification &dirSpec = construct.BeginDir();
+ mlir::SaveStateStack<OpenMPContextFrame> context{converter.getStateStack(),
+ eval, dirSpec.DirId()};
omp::List<omp::Clause> clauses = makeClauses(dirSpec.Clauses(), semaCtx);
lower::StatementContext stmtCtx;
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index df14d13b76a4bd..c09381b15be693 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -7502,8 +7502,7 @@ static void genMetadirective(lower::AbstractConverter &converter,
fir::FirOpBuilder &builder = converter.getFirOpBuilder();
llvm::SmallVector<llvm::omp::TraitProperty, 8> constructTraits;
- collectEnclosingConstructTraits(builder.getInsertionBlock()->getParentOp(),
- constructTraits);
+ collectEnclosingConstructTraits(converter, eval, constructTraits);
semantics::omp::OmpVariantMatchContext ompCtx =
makeVariantMatchContext(builder.getModule(), constructTraits);
@@ -7730,6 +7729,8 @@ static void genMetadirective(lower::AbstractConverter &converter,
if (marking == MetadirectiveLoopIVMarking::ThreadprivateIV)
TODO(variantLoc, "THREADPRIVATE loop iteration variable in "
"loop-associated METADIRECTIVE variant");
+ mlir::SaveStateStack<OpenMPContextFrame> context{
+ converter.getStateStack(), eval, spec->DirId()};
genOMPDispatch(converter, symTable, semaCtx, eval, variantLoc, queue,
queue.begin(), dsaGuard.getMarkedSymbols());
return;
@@ -7743,8 +7744,15 @@ static void genMetadirective(lower::AbstractConverter &converter,
TODO(variantLoc,
"METADIRECTIVE with both block- and loop-associated variants");
- genOMPDispatch(converter, symTable, semaCtx, eval, variantLoc, queue,
- queue.begin());
+ if (consumesBody) {
+ mlir::SaveStateStack<OpenMPContextFrame> context{
+ converter.getStateStack(), eval, spec->DirId()};
+ genOMPDispatch(converter, symTable, semaCtx, eval, variantLoc, queue,
+ queue.begin());
+ } else {
+ genOMPDispatch(converter, symTable, semaCtx, eval, variantLoc, queue,
+ queue.begin());
+ }
// A standalone variant (Association::None, e.g. barrier/taskwait/nothing)
// does not consume the metadirective's nested block, so lower it here.
if (!consumesBody && eval.hasNestedEvaluations())
diff --git a/flang/lib/Lower/OpenMP/Utils.cpp b/flang/lib/Lower/OpenMP/Utils.cpp
index b52c03ff63fed0..634ced155b0633 100644
--- a/flang/lib/Lower/OpenMP/Utils.cpp
+++ b/flang/lib/Lower/OpenMP/Utils.cpp
@@ -1473,27 +1473,51 @@ semantics::omp::OmpVariantMatchContext makeVariantMatchContext(
}
void collectEnclosingConstructTraits(
- mlir::Operation *op,
+ AbstractConverter &converter, const pft::Evaluation &evaluation,
llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits) {
- // Collect enclosing OpenMP operations so variants chosen by an outer
- // metadirective are part of this metadirective's context. For example, an
- // inner metadirective inside `target` and an outer-selected `parallel` must
- // 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))
- constructTraits.push_back(llvm::omp::TraitProperty::construct_for_for);
- if (mlir::isa<mlir::omp::ParallelOp>(op))
- constructTraits.push_back(
- llvm::omp::TraitProperty::construct_parallel_parallel);
- if (mlir::isa<mlir::omp::TeamsOp>(op))
- constructTraits.push_back(
- llvm::omp::TraitProperty::construct_teams_teams);
- if (mlir::isa<mlir::omp::TargetOp>(op))
- constructTraits.push_back(
- llvm::omp::TraitProperty::construct_target_target);
+ llvm::SmallVector<const OpenMPContextFrame *, 4> frames;
+ converter.getStateStack().stackWalk<OpenMPContextFrame>(
+ [&](OpenMPContextFrame &frame) {
+ frames.push_back(&frame);
+ return mlir::WalkResult::advance();
+ });
+ std::reverse(frames.begin(), frames.end());
+ llvm::SmallVector<bool, 4> usedFrames(frames.size(), false);
+
+ llvm::SmallVector<const pft::Evaluation *, 8> ancestors;
+ for (const pft::Evaluation *parent = evaluation.parentConstruct; parent;
+ parent = parent->parentConstruct) {
+ ancestors.push_back(parent);
+ }
+ std::reverse(ancestors.begin(), ancestors.end());
+
+ auto append = [&](llvm::omp::Directive directive) {
+ semantics::omp::AppendDirectiveContextTraits(directive, constructTraits);
+ };
+ for (const pft::Evaluation *ancestor : ancestors) {
+ const auto *omp = ancestor->getIf<parser::OpenMPConstruct>();
+ if (!omp)
+ continue;
+ llvm::omp::Directive directive{parser::omp::GetOmpDirectiveName(*omp).v};
+ if (directive != llvm::omp::Directive::OMPD_metadirective) {
+ append(directive);
+ continue;
+ }
+ for (auto [index, frame] : llvm::enumerate(frames)) {
+ if (&frame->evaluation == ancestor) {
+ append(frame->directive);
+ usedFrames[index] = true;
+ break;
+ }
+ }
}
- std::reverse(constructTraits.begin(), constructTraits.end());
+
+ // Active source contexts may not appear in the ancestor chain. This occurs
+ // for a loop-associated metadirective that owns a following sibling and for
+ // an atomic expression lowered from the atomic evaluation itself.
+ for (auto [index, frame] : llvm::enumerate(frames))
+ if (!usedFrames[index])
+ append(frame->directive);
}
const semantics::Symbol *
@@ -1537,9 +1561,8 @@ resolveDeclareVariantCallee(const semantics::Symbol &base,
}
llvm::SmallVector<llvm::omp::TraitProperty, 8> constructTraits;
- collectEnclosingConstructTraits(
- converter.getFirOpBuilder().getInsertionBlock()->getParentOp(),
- constructTraits);
+ collectEnclosingConstructTraits(converter, converter.getCurrentEvaluation(),
+ constructTraits);
semantics::omp::OmpVariantMatchContext ompCtx =
makeVariantMatchContext(converter.getModuleOp(), constructTraits);
diff --git a/flang/lib/Lower/OpenMP/Utils.h b/flang/lib/Lower/OpenMP/Utils.h
index 7fde1825b006a2..7b89220f432f65 100644
--- a/flang/lib/Lower/OpenMP/Utils.h
+++ b/flang/lib/Lower/OpenMP/Utils.h
@@ -15,6 +15,7 @@
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "mlir/IR/Location.h"
#include "mlir/IR/Value.h"
+#include "mlir/Support/StateStack.h"
#include "llvm/Frontend/OpenMP/OMPContext.h"
#include "llvm/Support/CommandLine.h"
#include <cstdint>
@@ -51,6 +52,19 @@ class AbstractConverter;
namespace omp {
+class OpenMPContextFrame
+ : public mlir::StateStackFrameBase<OpenMPContextFrame> {
+public:
+ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(OpenMPContextFrame)
+
+ OpenMPContextFrame(const pft::Evaluation &evaluation,
+ llvm::omp::Directive directive)
+ : evaluation{evaluation}, directive{directive} {}
+
+ const pft::Evaluation &evaluation;
+ llvm::omp::Directive directive;
+};
+
struct DeclareTargetCaptureInfo {
mlir::omp::DeclareTargetCaptureClause clause;
bool automap = false;
@@ -262,12 +276,11 @@ std::optional<llvm::SmallVector<mlir::Value>> getIteratorElementIndices(
Fortran::lower::AbstractConverter &converter, const omp::Object &object,
Fortran::lower::StatementContext &stmtCtx, mlir::Location loc);
-/// 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`.
+/// Collect the source OpenMP constructs enclosing \p evaluation in
+/// outermost-to-innermost order. Active metadirective replacements substitute
+/// for their METADIRECTIVE source constructs.
void collectEnclosingConstructTraits(
- mlir::Operation *op,
+ AbstractConverter &converter, const pft::Evaluation &evaluation,
llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits);
/// Return true when \p module is being compiled for an AMDGPU device or all of
diff --git a/flang/lib/Semantics/openmp-utils.cpp b/flang/lib/Semantics/openmp-utils.cpp
index a341cf4fa17065..6e311cd2e38f4f 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);
}
}
}
@@ -2453,6 +2452,36 @@ static void AppendConstructTraitsForDirective(
add(llvm::omp::TraitProperty::construct_dispatch_dispatch);
}
+void AppendDirectiveContextTraits(llvm::omp::Directive directive,
+ llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits) {
+ using llvm::omp::Directive;
+ using llvm::omp::TraitProperty;
+
+ for (Directive leaf : llvm::omp::getLeafConstructsOrSelf(directive)) {
+ if (leaf == Directive::OMPD_nothing ||
+ leaf == Directive::OMPD_metadirective ||
+ leaf == Directive::OMPD_unknown || leaf == Directive::OMPD_section ||
+ leaf == Directive::OMPD_dispatch ||
+ llvm::omp::getDirectiveCategory(leaf) ==
+ llvm::omp::Category::Informational) {
+ continue;
+ }
+
+ llvm::omp::VariantMatchInfo vmi;
+ AppendConstructTraitsForDirective(leaf, vmi);
+ if (vmi.RequiredTraits.test(
+ static_cast<unsigned>(TraitProperty::construct_target_target))) {
+ constructTraits.clear();
+ }
+ if (vmi.ConstructTraits.empty()) {
+ constructTraits.push_back(TraitProperty::invalid);
+ } else {
+ constructTraits.append(
+ vmi.ConstructTraits.begin(), vmi.ConstructTraits.end());
+ }
+ }
+}
+
static void AddTraitPropertiesFromSelector(llvm::omp::TraitSet set,
const parser::OmpTraitSelector &selector, llvm::omp::VariantMatchInfo &vmi,
SemanticsContext &semaCtx,
@@ -2480,17 +2509,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 +2535,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 +2637,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 +2650,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)) {
@@ -2629,33 +2667,32 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet(
auto addConditionTraitForRanking =
[&](llvm::omp::VariantMatchInfo &rankingVMI) {
rankingVMI.addTrait(hasMatchNone
- ? dynamicConditionTrait
+ ? llvm::omp::TraitProperty::user_condition_false
: 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});
- }
- result.candidates.push_back({spec, std::move(staticVMI), isExplicit});
+ // Represent both outcomes. Keeping the false condition in the
+ // unguarded candidate preserves selector identity for subset
+ // comparisons; MATCH_ANY scoring omits its inactive score.
+ llvm::omp::VariantMatchInfo conditionTrueVMI{staticVMI};
+ addConditionTraitForRanking(conditionTrueVMI);
+ llvm::omp::VariantMatchInfo conditionFalseVMI{staticVMI};
+ conditionFalseVMI.addTrait(
+ llvm::omp::TraitProperty::user_condition_false, conditionSource,
+ conditionScorePtr);
+ result.candidates.push_back({spec, std::move(conditionTrueVMI),
+ isExplicit, dynamicCondition});
+ result.candidates.push_back(
+ {spec, std::move(conditionFalseVMI), 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 +2926,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 +2987,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 ca72b5c8e47af8..90fe7250b9bd64 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/declare-variant-source-context.f90 b/flang/test/Lower/OpenMP/declare-variant-source-context.f90
new file mode 100644
index 00000000000000..420c649c44456e
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-variant-source-context.f90
@@ -0,0 +1,113 @@
+! RUN: %if x86-registered-target %{ %flang_fc1 -fopenmp \
+! RUN: -fopenmp-version=52 -triple x86_64-unknown-linux-gnu \
+! RUN: -emit-hlfir %s -o - | FileCheck %s %}
+
+module source_context
+contains
+ subroutine depth_base
+ !$omp declare variant(depth_vendor) &
+ !$omp& match(implementation={vendor(score(3): llvm)})
+ !$omp declare variant(depth_arch) match(device={arch(x86_64)})
+ end subroutine
+ subroutine depth_vendor
+ end subroutine
+ subroutine depth_arch
+ end subroutine
+
+ subroutine order_base
+ !$omp declare variant(order_vendor) &
+ !$omp& match(implementation={vendor(score(3): llvm)})
+ !$omp declare variant(order_parallel) match(construct={parallel})
+ end subroutine
+ subroutine order_vendor
+ end subroutine
+ subroutine order_parallel
+ end subroutine
+
+ integer function value_base()
+ !$omp declare variant(value_vendor) &
+ !$omp& match(implementation={vendor(score(3): llvm)})
+ !$omp declare variant(value_arch) match(device={arch(x86_64)})
+ value_base = 0
+ end function
+ integer function value_vendor()
+ value_vendor = 1
+ end function
+ integer function value_arch()
+ value_arch = 2
+ end function
+
+! CHECK-LABEL: func.func @_QMsource_contextPtile_context(
+! CHECK: fir.call @_QMsource_contextPdepth_arch()
+! CHECK-NOT: fir.call @_QMsource_contextPdepth_vendor
+! CHECK: return
+ subroutine tile_context(n)
+ integer :: n, i
+ !$omp tile sizes(2)
+ do i = 1, n
+ call depth_base()
+ end do
+ end subroutine
+
+! CHECK-LABEL: func.func @_QMsource_contextPunroll_context(
+! CHECK: fir.call @_QMsource_contextPdepth_arch()
+! CHECK-NOT: fir.call @_QMsource_contextPdepth_vendor
+! CHECK: return
+ subroutine unroll_context(n)
+ integer :: n, i
+ !$omp unroll partial(2)
+ do i = 1, n
+ call depth_base()
+ end do
+ end subroutine
+
+! CHECK-LABEL: func.func @_QMsource_contextPfuse_context(
+! CHECK-NOT: fir.call @_QMsource_contextPdepth_vendor
+! CHECK-COUNT-2: fir.call @_QMsource_contextPdepth_arch()
+! CHECK-NOT: fir.call @_QMsource_contextPdepth_vendor
+! CHECK: return
+ subroutine fuse_context(n)
+ integer :: n, i, j
+ !$omp fuse
+ do i = 1, n
+ call depth_base()
+ end do
+ do j = 1, n
+ call depth_base()
+ end do
+ !$omp end fuse
+ end subroutine
+
+! CHECK-LABEL: func.func @_QMsource_contextPatomic_context(
+! CHECK: fir.call @_QMsource_contextPvalue_arch()
+! CHECK-NOT: fir.call @_QMsource_contextPvalue_vendor
+! CHECK: omp.atomic.update
+! CHECK: return
+ subroutine atomic_context(x)
+ integer :: x
+ !$omp atomic update
+ x = x + value_base()
+ end subroutine
+
+! CHECK-LABEL: func.func @_QMsource_contextPassume_context()
+! CHECK: fir.call @_QMsource_contextPdepth_vendor()
+! CHECK-NOT: fir.call @_QMsource_contextPdepth_arch
+! CHECK: return
+ subroutine assume_context
+ !$omp assume holds(.true.)
+ call depth_base()
+ !$omp end assume
+ end subroutine
+
+! CHECK-LABEL: func.func @_QMsource_contextPcombined_context(
+! CHECK: fir.call @_QMsource_contextPorder_parallel()
+! CHECK-NOT: fir.call @_QMsource_contextPorder_vendor
+! CHECK: return
+ subroutine combined_context(n)
+ integer :: n, i
+ !$omp teams distribute parallel do
+ do i = 1, n
+ call order_base()
+ end do
+ end subroutine
+end module
diff --git a/flang/test/Lower/OpenMP/metadirective-device-kind.f90 b/flang/test/Lower/OpenMP/metadirective-device-kind.f90
index ae7613d40c5d90..d96be7739ac98c 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 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/flang/test/Lower/OpenMP/metadirective-source-context.f90 b/flang/test/Lower/OpenMP/metadirective-source-context.f90
new file mode 100644
index 00000000000000..9472dd6e92afc5
--- /dev/null
+++ b/flang/test/Lower/OpenMP/metadirective-source-context.f90
@@ -0,0 +1,68 @@
+! RUN: %if x86-registered-target %{ %flang_fc1 -fopenmp \
+! RUN: -fopenmp-version=52 -triple x86_64-unknown-linux-gnu \
+! RUN: -emit-hlfir %s -o - | FileCheck %s %}
+
+! Executable loop transformations contribute to the construct context.
+! CHECK-LABEL: func.func @_QPtile_context(
+! CHECK: omp.barrier
+! CHECK-NOT: omp.taskyield
+! CHECK: return
+subroutine tile_context(n, a)
+ integer :: n, a(n, n), i, j
+ !$omp tile sizes(2)
+ do i = 1, n
+ !$omp metadirective &
+ !$omp& when(implementation={vendor(score(3): llvm)}: taskyield) &
+ !$omp& when(device={arch(x86_64)}: barrier)
+ do j = 1, n
+ a(j, i) = j
+ end do
+ end do
+end subroutine
+
+! CHECK-LABEL: func.func @_QPunroll_context(
+! CHECK: omp.barrier
+! CHECK-NOT: omp.taskyield
+! CHECK: return
+subroutine unroll_context(n, a)
+ integer :: n, a(n), i
+ !$omp unroll partial(2)
+ do i = 1, n
+ !$omp metadirective &
+ !$omp& when(implementation={vendor(score(3): llvm)}: taskyield) &
+ !$omp& when(device={arch(x86_64)}: barrier)
+ a(i) = i
+ end do
+end subroutine
+
+! Informational directives do not contribute to the construct context.
+! CHECK-LABEL: func.func @_QPassume_context()
+! CHECK: omp.taskyield
+! CHECK-NOT: omp.barrier
+! CHECK: return
+subroutine assume_context()
+ !$omp assume holds(.true.)
+ !$omp metadirective &
+ !$omp& when(implementation={vendor(score(3): llvm)}: taskyield) &
+ !$omp& when(device={arch(x86_64)}: barrier)
+ !$omp end assume
+end subroutine
+
+! Combined and composite constructs retain their source nesting order.
+! CHECK-LABEL: func.func @_QPcombined_context(
+! CHECK: omp.teams
+! CHECK: omp.parallel
+! CHECK: omp.distribute
+! CHECK: omp.wsloop
+! CHECK: omp.taskyield
+! CHECK-NOT: omp.barrier
+! CHECK: return
+subroutine combined_context(n)
+ integer :: n, i
+ !$omp teams distribute parallel do
+ do i = 1, n
+ !$omp metadirective &
+ !$omp& when(implementation={vendor(score(3): llvm)}: barrier) &
+ !$omp& when(construct={parallel}: taskyield)
+ end do
+end subroutine
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 00000000000000..d7963bf7fe4ece
--- /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-static.f90 b/flang/test/Lower/OpenMP/metadirective-user-static.f90
new file mode 100644
index 00000000000000..340b7c25293c87
--- /dev/null
+++ b/flang/test/Lower/OpenMP/metadirective-user-static.f90
@@ -0,0 +1,81 @@
+! Test lowering of OpenMP metadirective with static user={condition()}
+! selectors.
+
+! RUN: %flang_fc1 -fopenmp -emit-hlfir -fopenmp-version=50 %s -o - | \
+! RUN: FileCheck %s
+! RUN: %flang_fc1 -fopenmp -emit-hlfir -fopenmp-version=51 %s -o - | \
+! RUN: FileCheck %s
+! RUN: %flang_fc1 -fopenmp -emit-hlfir -fopenmp-version=52 %s -o - | \
+! RUN: FileCheck %s
+
+! 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
+
+! CHECK-LABEL: func.func @_QPtest_condition_true()
+! CHECK: omp.taskyield
+! CHECK-NOT: fir.if
+! CHECK: return
+subroutine test_condition_true()
+ !$omp metadirective &
+ !$omp & when(user={condition(.true.)}: taskyield)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_condition_false()
+! CHECK-NOT: omp.taskwait
+! CHECK-NOT: fir.if
+! CHECK: return
+subroutine test_condition_false()
+ !$omp metadirective &
+ !$omp & when(user={condition(.false.)}: taskwait)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_condition_score()
+! CHECK-NOT: omp.taskyield
+! CHECK: omp.taskwait
+! CHECK: return
+subroutine test_condition_score()
+ !$omp metadirective &
+ !$omp & when(user={condition(.true.)}: taskyield) &
+ !$omp & when(user={condition(score(2): .true.)}: taskwait)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_begin_condition_true()
+! CHECK: omp.parallel
+! CHECK: omp.terminator
+! CHECK-NOT: fir.if
+! CHECK: return
+subroutine test_begin_condition_true()
+ integer :: x
+ x = 0
+ !$omp begin metadirective &
+ !$omp & when(user={condition(.true.)}: parallel)
+ x = 1
+ !$omp end metadirective
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_begin_condition_false()
+! CHECK-NOT: omp.parallel
+! CHECK-NOT: fir.if
+! CHECK: return
+subroutine test_begin_condition_false()
+ integer :: x
+ x = 0
+ !$omp begin metadirective &
+ !$omp & when(user={condition(.false.)}: parallel)
+ x = 1
+ !$omp end metadirective
+end subroutine
diff --git a/flang/test/Lower/OpenMP/metadirective-user.f90 b/flang/test/Lower/OpenMP/metadirective-user.f90
index cdfbddd4151a03..23c92954e7e75d 100644
--- a/flang/test/Lower/OpenMP/metadirective-user.f90
+++ b/flang/test/Lower/OpenMP/metadirective-user.f90
@@ -1,93 +1,169 @@
-! Test lowering of OpenMP metadirective with user={condition()} selectors.
+! Test lowering of OpenMP metadirective with dynamic user={condition()}
+! selectors.
-! RUN: %flang_fc1 -fopenmp -emit-hlfir -fopenmp-version=50 %s -o - | FileCheck %s
! 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
!===----------------------------------------------------------------------===!
-! Static (constant-folded) user conditions
+! 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
+
+! Scored implicit NOTHING competes with explicit replacements by score.
!===----------------------------------------------------------------------===!
-! CHECK-LABEL: func.func @_QPtest_condition_true()
-! CHECK: omp.taskyield
-! CHECK-NOT: fir.if
-! CHECK: return
-subroutine test_condition_true()
+! 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(.true.)}: taskyield) &
-#ifdef OMP_52
- !$omp & otherwise(nothing)
-#else
- !$omp & default(nothing)
-#endif
+ !$omp& when(user={condition(score(10): flag)}:) &
+ !$omp& when(user={condition(score(5): .true.)}: barrier)
end subroutine
-! CHECK-LABEL: func.func @_QPtest_condition_false()
-! CHECK-NOT: omp.taskwait
-! CHECK-NOT: fir.if
-! CHECK: return
-subroutine test_condition_false()
+! 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(.false.)}: taskwait) &
-#ifdef OMP_52
- !$omp & otherwise(nothing)
-#else
- !$omp & default(nothing)
-#endif
+ !$omp& when(user={condition(score(10): flag)}: nothing) &
+ !$omp& when(user={condition(score(5): .true.)}: barrier)
end subroutine
-! CHECK-LABEL: func.func @_QPtest_condition_score()
-! CHECK-NOT: omp.taskyield
-! CHECK: omp.taskwait
-! CHECK: return
-subroutine test_condition_score()
+! 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(.true.)}: taskyield) &
- !$omp & when(user={condition(score(2): .true.)}: taskwait) &
-#ifdef OMP_52
- !$omp & otherwise(nothing)
-#else
- !$omp & default(nothing)
-#endif
+ !$omp& when(user={condition(score(5): flag)}:) &
+ !$omp& when(user={condition(score(5): .true.)}: barrier)
end subroutine
-! CHECK-LABEL: func.func @_QPtest_begin_condition_true()
-! CHECK: omp.parallel
-! CHECK: omp.terminator
-! CHECK-NOT: fir.if
-! CHECK: return
-subroutine test_begin_condition_true()
- integer :: x
- x = 0
-#ifdef OMP_52
- !$omp begin metadirective &
- !$omp & when(user={condition(.true.)}: parallel) &
- !$omp & otherwise(nothing)
-#else
- !$omp begin metadirective &
- !$omp & when(user={condition(.true.)}: parallel)
-#endif
- x = 1
- !$omp end metadirective
+! 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
-! CHECK-LABEL: func.func @_QPtest_begin_condition_false()
-! CHECK-NOT: omp.parallel
-! CHECK-NOT: fir.if
-! CHECK: return
-subroutine test_begin_condition_false()
- integer :: x
- x = 0
-#ifdef OMP_52
- !$omp begin metadirective &
- !$omp & when(user={condition(.false.)}: parallel) &
- !$omp & otherwise(nothing)
-#else
- !$omp begin metadirective &
- !$omp & when(user={condition(.false.)}: parallel)
-#endif
- x = 1
- !$omp end metadirective
+! 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
!===----------------------------------------------------------------------===!
@@ -331,8 +407,21 @@ subroutine test_dynamic_user_score_order(low, high)
!$omp & when(device={kind(host)}: taskwait)
end subroutine
-! Under extension(match_none), the runtime condition selects the variant when it
-! is false. Its score is still available for ranking that false path.
+! A statically false MATCH_NONE condition is the selector that makes the
+! candidate applicable, so its score participates in ranking.
+! CHECK-LABEL: func.func @_QPtest_static_user_score_match_none()
+! CHECK-NOT: fir.if
+! CHECK: omp.barrier
+! CHECK-NOT: omp.taskwait
+! CHECK: return
+subroutine test_static_user_score_match_none()
+ !$omp metadirective &
+ !$omp & when(implementation={extension(match_none)}, &
+ !$omp & user={condition(score(100): .false.)}: barrier) &
+ !$omp & when(user={condition(score(10): .true.)}: taskwait)
+end subroutine
+
+! The dynamic form must use the same score on its false path.
! CHECK-LABEL: func.func @_QPtest_dynamic_user_score_match_none(
! CHECK-SAME: %[[ARG0:.*]]: !fir.ref<!fir.logical<4>>
! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[ARG0]]
@@ -349,12 +438,9 @@ subroutine test_dynamic_user_score_order(low, high)
subroutine test_dynamic_user_score_match_none(flag)
logical, intent(in) :: flag
!$omp metadirective &
- !$omp & when(implementation={extension(match_none)}, user={condition(score(5): flag)}: barrier) &
-#ifdef OMP_52
- !$omp & otherwise(taskwait)
-#else
- !$omp & default(taskwait)
-#endif
+ !$omp & when(implementation={extension(match_none)}, &
+ !$omp & user={condition(score(100): flag)}: barrier) &
+ !$omp & when(user={condition(score(10): .true.)}: taskwait)
end subroutine
! Under extension(match_any), a dynamic condition can be the selector that makes
@@ -425,13 +511,61 @@ 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
+! A statically false condition follows the dynamic false-path rule: its score
+! does not contribute when another selector satisfies MATCH_ANY.
+! CHECK-LABEL: func.func @_QPtest_static_user_match_any_false_score()
+! CHECK-NOT: fir.if
+! CHECK-NOT: omp.barrier
+! CHECK: omp.taskwait
+! CHECK: return
+subroutine test_static_user_match_any_false_score()
+ !$omp metadirective &
+ !$omp& when(implementation={extension(match_any), vendor(llvm)}, &
+ !$omp& user={condition(score(100): .false.)}: barrier) &
+ !$omp& when(user={condition(score(10): .true.)}: taskwait) &
+#ifdef OMP_52
+ !$omp& otherwise(nothing)
+#else
+ !$omp& default(nothing)
+#endif
+end subroutine
+
+! A false user condition remains part of the selector for specificity. The
+! more-specific MATCH_ANY candidate therefore wins over the vendor-only
+! candidate in both the static and dynamic forms.
+! CHECK-LABEL: func.func @_QPtest_static_user_match_any_false_subset()
+! CHECK-NOT: fir.if
+! CHECK: omp.barrier
+! CHECK-NOT: omp.taskwait
+! CHECK: return
+subroutine test_static_user_match_any_false_subset()
+ !$omp metadirective &
+ !$omp& when(implementation={extension(match_any), vendor(llvm)}, &
+ !$omp& user={condition(.false.)}: barrier) &
+ !$omp& when(implementation={vendor(llvm)}: taskwait)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_dynamic_user_match_any_false_subset(
+! CHECK-NOT: fir.if
+! CHECK: omp.barrier
+! CHECK-NOT: omp.taskwait
+! CHECK: return
+subroutine test_dynamic_user_match_any_false_subset(flag)
+ logical, intent(in) :: flag
+ !$omp metadirective &
+ !$omp& when(implementation={extension(match_any), vendor(llvm)}, &
+ !$omp& user={condition(flag)}: barrier) &
+ !$omp& when(implementation={vendor(llvm)}: taskwait)
+end subroutine
+
+! 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 00000000000000..e51df9bc37746a
--- /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 7849d326659944..f6f3f3a22eeff4 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,36 @@ 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.
+/// Variant match information describes the required property traits, how they
+/// are scored, and the required construct nesting.
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 +169,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 +209,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 f2cdd9bbaf5e45..69d8c94cf98be5 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,120 @@ 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();
+ bool HasMatchAny = VMI.RequiredTraits.test(
+ unsigned(TraitProperty::implementation_extension_match_any));
+ 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. A false condition contributes no
+ // score when match_any is satisfied by another selector.
+ if (getOpenMPContextTraitSetForProperty(Property) == TraitSet::construct ||
+ Property == TraitProperty::device_kind_any ||
+ Property == TraitProperty::target_device_kind_any ||
+ (Property == TraitProperty::user_condition_false && HasMatchAny))
+ 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 f9683ae56e933d..cf8e68f0a33a5b 100644
--- a/llvm/unittests/Frontend/OpenMPContextTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPContextTest.cpp
@@ -314,7 +314,536 @@ 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 MATCH_ANY beats the
+ // explicitly scored candidate's 1 + 2. Retaining only either construct
+ // match would not select MATCH_ANY.
+ 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 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, ScoringMatchAnyFalseCondition) {
+ OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1);
+ APInt HighScore(64, 100), LowScore(64, 10);
+
+ VariantMatchInfo MatchAnyFalse;
+ MatchAnyFalse.addTrait(TraitProperty::implementation_vendor_llvm, "");
+ MatchAnyFalse.addTrait(TraitProperty::implementation_extension_match_any, "");
+ MatchAnyFalse.addTrait(TraitProperty::user_condition_false, "condition",
+ &HighScore);
+
+ VariantMatchInfo ScoredTrue;
+ ScoredTrue.addTrait(TraitProperty::user_condition_true, "other condition",
+ &LowScore);
+
+ // A matching static trait keeps MATCH_ANY applicable, but a false condition
+ // does not contribute its score.
+ SmallVector<VariantMatchInfo, 2> Candidates{MatchAnyFalse, ScoredTrue};
+ EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 1);
+
+ VariantMatchInfo MatchAnyTrue;
+ MatchAnyTrue.addTrait(TraitProperty::implementation_vendor_llvm, "");
+ MatchAnyTrue.addTrait(TraitProperty::implementation_extension_match_any, "");
+ MatchAnyTrue.addTrait(TraitProperty::user_condition_true, "condition",
+ &HighScore);
+ Candidates[0] = MatchAnyTrue;
+ EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0);
+
+ VariantMatchInfo MatchNoneFalse;
+ MatchNoneFalse.addTrait(TraitProperty::implementation_extension_match_none,
+ "");
+ MatchNoneFalse.addTrait(TraitProperty::user_condition_false, "condition",
+ &HighScore);
+ // A false condition is the selector that makes MATCH_NONE applicable, so
+ // its score still participates in ranking.
+ Candidates[0] = MatchNoneFalse;
+ EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0);
+}
+
+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);
+
+ APInt HighScore(64, 100);
+ First = VariantMatchInfo();
+ First.addTrait(TraitProperty::implementation_extension_match_none, "");
+ First.addUnknownTrait(TraitSelector::implementation_vendor, "same",
+ &HighScore);
+ Second = VariantMatchInfo();
+ Second.addTrait(TraitProperty::implementation_extension_match_none, "");
+ Second.addUnknownTrait(TraitSelector::device_kind, "same");
+ Second.addUnknownTrait(TraitSelector::device_arch, "extra");
+ Candidates = {First, Second};
+ // The same property name under different selectors does not form a subset.
+ EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0);
+}
+
+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, IncompatibleSupersetDoesNotZeroScore) {
+ OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1);
+ APInt HighScore(64, 100), LowScore(64, 5);
+ VariantMatchInfo Subset, IncompatibleSuperset, Other;
+ Subset.addTrait(TraitProperty::implementation_vendor_llvm, "", &HighScore);
+ IncompatibleSuperset.addTrait(TraitProperty::implementation_vendor_llvm, "");
+ IncompatibleSuperset.addTrait(TraitProperty::user_condition_false, "false");
+ Other.addTrait(TraitProperty::user_condition_true, "true", &LowScore);
+
+ // Only compatible selectors participate in strict-subset zeroing.
+ SmallVector<VariantMatchInfo, 3> Candidates{Subset, IncompatibleSuperset,
+ Other};
+ 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");
+ High.addTrait(TraitProperty::target_device_isa___ANY, "avx");
+ Low.addTrait(TraitProperty::implementation_vendor_llvm, "", &LowScore);
+ Low.addTrait(TraitProperty::device_isa___ANY, "avx");
+ Low.addTrait(TraitProperty::target_device_isa___ANY, "sse2");
+ 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 and their selectors are identical, HIGH is a
+ // strict subset of LOW.
+ Candidates[0].ISATraits[0].Name = "sse2";
+ Candidates[0].ISATraits[1].Name = "avx";
+ 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
>From e5d4cbc278813eff199685e38e612e8c7722d9e9 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Thu, 24 Sep 2026 00:05:55 -0500
Subject: [PATCH 2/2] Fix variant matching in expressions
- Fix Flang crashes when selecting variants in specification expressions.
- Make variant selection in combined-loop bounds agree with explicit nesting.
- Fix Clang condition matching across macro redefinitions and parentheses.
- Strengthen tests for selected TARGET contexts and add bug regressions.
---
clang/lib/Parse/ParseOpenMP.cpp | 13 +-
.../declare_variant_condition_identity_pch.c | 20 ++-
clang/test/OpenMP/declare_variant_scoring.c | 44 +++++++
flang/include/flang/Lower/AbstractConverter.h | 4 +-
flang/lib/Lower/Bridge.cpp | 6 +-
flang/lib/Lower/OpenMP/OpenMP.cpp | 17 ++-
flang/lib/Lower/OpenMP/Utils.cpp | 37 ++++--
flang/lib/Lower/OpenMP/Utils.h | 12 +-
.../OpenMP/declare-variant-loop-bounds.f90 | 120 ++++++++++++++++++
.../OpenMP/declare-variant-spec-expr.f90 | 56 ++++++++
.../OpenMP/metadirective-target-boundary.f90 | 34 +++--
11 files changed, 323 insertions(+), 40 deletions(-)
create mode 100644 flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
create mode 100644 flang/test/Lower/OpenMP/declare-variant-spec-expr.f90
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 6f558df8e8b851..c6a358f8af823e 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -28,6 +28,7 @@
#include "llvm/Frontend/OpenMP/DirectiveNameParser.h"
#include "llvm/Frontend/OpenMP/OMPAssume.h"
#include "llvm/Frontend/OpenMP/OMPContext.h"
+#include "llvm/Support/raw_ostream.h"
#include <climits>
#include <optional>
@@ -1150,12 +1151,16 @@ void Parser::parseOMPContextSelector(
if (!Condition.isUsable())
return FinishSelector();
TISelector.ScoreOrCondition = Condition.get();
- StringRef ConditionText = Lexer::getSourceText(
- CharSourceRange::getTokenRange(Condition.get()->getSourceRange()),
- PP.getSourceManager(), PP.getLangOpts());
+ // Compare expanded expressions, since a macro can be redefined between
+ // selectors. Ignore outer parentheses when recording condition identity.
+ ASTContext &Context = Actions.getASTContext();
+ std::string ConditionText;
+ llvm::raw_string_ostream OS(ConditionText);
+ Condition.get()->IgnoreParenImpCasts()->printPretty(
+ OS, nullptr, Context.getPrintingPolicy());
TISelector.Properties.push_back(
{TraitProperty::user_condition_unknown,
- ConditionText.empty() ? "<condition>" : ConditionText});
+ StringRef(ConditionText).copy(Context.getAllocator())});
return;
}
diff --git a/clang/test/OpenMP/declare_variant_condition_identity_pch.c b/clang/test/OpenMP/declare_variant_condition_identity_pch.c
index 3b0acd9f8f1f6b..166612af3a9335 100644
--- a/clang/test/OpenMP/declare_variant_condition_identity_pch.c
+++ b/clang/test/OpenMP/declare_variant_condition_identity_pch.c
@@ -12,7 +12,7 @@
// RUN: -triple x86_64-unknown-linux -x c++ %t/use.c \
// RUN: -include-pch %t/conditions-cxx.pch -emit-llvm -o - | FileCheck %s
-// Verify that serialization preserves the source identity of folded user
+// Verify that serialization preserves the expanded identity of folded user
// conditions. Distinct conditions do not form a subset relationship, while
// identical conditions do.
@@ -29,19 +29,29 @@ extern "C" {
void condition_high_variant(void);
void condition_low_variant(void);
+#define COND 1
#pragma omp declare variant(condition_high_variant) \
- match(implementation = {vendor(score(100) : llvm)}, user = {condition(1)})
+ match(implementation = {vendor(score(100) : llvm)}, \
+ user = {condition(COND)})
+#undef COND
+#define COND 2
#pragma omp declare variant(condition_low_variant) \
match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
- user = {condition(2)})
+ user = {condition(COND)})
void distinct_condition_base(void);
+#undef COND
+#define FIRST_COND 1
+#define SECOND_COND (1)
#pragma omp declare variant(condition_high_variant) \
- match(implementation = {vendor(score(100) : llvm)}, user = {condition(1)})
+ match(implementation = {vendor(score(100) : llvm)}, \
+ user = {condition(FIRST_COND)})
#pragma omp declare variant(condition_low_variant) \
match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
- user = {condition(1)})
+ user = {condition(SECOND_COND)})
void identical_condition_base(void);
+#undef FIRST_COND
+#undef SECOND_COND
#ifdef __cplusplus
}
diff --git a/clang/test/OpenMP/declare_variant_scoring.c b/clang/test/OpenMP/declare_variant_scoring.c
index c924dbaf346246..ea99df0fee2c82 100644
--- a/clang/test/OpenMP/declare_variant_scoring.c
+++ b/clang/test/OpenMP/declare_variant_scoring.c
@@ -4,6 +4,10 @@
// 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
+// RUN: %clang_cc1 -E -fopenmp -fopenmp-version=52 %s -o %t.i
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -target-feature +avx \
+// RUN: -emit-llvm %t.i -o - | FileCheck %s
// expected-no-diagnostics
#ifdef __cplusplus
@@ -61,6 +65,32 @@ void distinct_condition_base(void);
user = {condition(1)})
void identical_condition_base(void);
+// Macro spelling must not hide different expanded conditions.
+#define COND 1
+#pragma omp declare variant(condition_high_variant) \
+ match(implementation = {vendor(score(100) : llvm)}, \
+ user = {condition(COND)})
+#undef COND
+#define COND 2
+#pragma omp declare variant(condition_low_variant) \
+ match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
+ user = {condition(COND)})
+void redefined_condition_base(void);
+#undef COND
+
+// Different macro names and outer parentheses preserve condition identity.
+#define FIRST_COND 1
+#define SECOND_COND (1)
+#pragma omp declare variant(condition_high_variant) \
+ match(implementation = {vendor(score(100) : llvm)}, \
+ user = {condition(FIRST_COND)})
+#pragma omp declare variant(condition_low_variant) \
+ match(implementation = {vendor(score(1) : llvm)}, device = {kind(cpu)}, \
+ user = {condition(SECOND_COND)})
+void equivalent_condition_base(void);
+#undef FIRST_COND
+#undef SECOND_COND
+
void subset_variant(void);
void superset_variant(void);
#pragma omp declare variant(subset_variant) \
@@ -230,6 +260,20 @@ void identical_conditions(void) { identical_condition_base(); }
// CHECK: call void @condition_low_variant()
// CHECK: ret void
+void redefined_conditions(void) { redefined_condition_base(); }
+// CHECK-LABEL: define{{.*}} void @redefined_conditions
+// CHECK-NOT: call void @condition_low_variant()
+// CHECK: call void @condition_high_variant()
+// CHECK-NOT: call void @condition_low_variant()
+// CHECK: ret void
+
+void equivalent_conditions(void) { equivalent_condition_base(); }
+// CHECK-LABEL: define{{.*}} void @equivalent_conditions
+// CHECK-NOT: call void @condition_high_variant()
+// CHECK: call void @condition_low_variant()
+// CHECK-NOT: call void @condition_high_variant()
+// CHECK: ret void
+
// 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(); }
diff --git a/flang/include/flang/Lower/AbstractConverter.h b/flang/include/flang/Lower/AbstractConverter.h
index eb17d03d458535..e08e37ef59d5d1 100644
--- a/flang/include/flang/Lower/AbstractConverter.h
+++ b/flang/include/flang/Lower/AbstractConverter.h
@@ -337,8 +337,8 @@ class AbstractConverter {
/// Get the converter's current scope
virtual const Fortran::semantics::Scope &getCurrentScope() = 0;
- /// Get the PFT evaluation currently being lowered.
- virtual pft::Evaluation &getCurrentEvaluation() = 0;
+ /// Get the current PFT evaluation, or null when lowering a function prologue.
+ virtual pft::Evaluation *getCurrentEvaluation() = 0;
//===--------------------------------------------------------------------===//
// FIR/MLIR
diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp
index f1c127e01513ae..029b324c5cefa3 100644
--- a/flang/lib/Lower/Bridge.cpp
+++ b/flang/lib/Lower/Bridge.cpp
@@ -1229,8 +1229,8 @@ class FirConverter : public Fortran::lower::AbstractConverter {
return bridge.getSemanticsContext().FindScope(currentPosition);
}
- Fortran::lower::pft::Evaluation &getCurrentEvaluation() override final {
- return getEval();
+ Fortran::lower::pft::Evaluation *getCurrentEvaluation() override final {
+ return evalPtr;
}
fir::FirOpBuilder &getFirOpBuilder() override final {
@@ -6255,6 +6255,8 @@ class FirConverter : public Fortran::lower::AbstractConverter {
/// Start translation of a function.
void startNewFunction(Fortran::lower::pft::FunctionLikeUnit &funit) {
assert(!builder && "expected nullptr");
+ // Specification expressions have no current executable evaluation.
+ evalPtr = nullptr;
bridge.fctCtx().pushScope();
bridge.cudaCleanupCtx().pushScope();
bridge.openAccCtx().pushScope();
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index c09381b15be693..190349bcb5c061 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -2017,6 +2017,9 @@ static void createBodyOfOp(mlir::Operation &op, const OpWithBodyGenInfo &info,
return {};
}();
+ mlir::SaveStateStack<OpenMPContextFrame> context{
+ info.converter.getStateStack(), info.eval, info.dir, /*isPartial=*/true};
+
// Mark the earliest insertion point.
mlir::Operation *marker = insertMarker(firOpBuilder);
@@ -2167,6 +2170,9 @@ static void genBodyOfTargetDataOp(
fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
genEntryBlock(firOpBuilder, args.asEntryBlockArgs(), dataOp.getRegion());
+ mlir::SaveStateStack<OpenMPContextFrame> context{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_target_data,
+ /*isPartial=*/true};
bindEntryBlockArgs(converter, dataOp, args);
auto argIface = llvm::cast<mlir::omp::BlockArgOpenMPOpInterface>(*dataOp);
llvm::SmallVector<const semantics::Symbol *> sourceUseDeviceAddrSyms{
@@ -2253,6 +2259,9 @@ static void genBodyOfTargetOp(
mlir::Region ®ion = targetOp.getRegion();
genEntryBlock(firOpBuilder, args.asEntryBlockArgs(), region);
+ mlir::SaveStateStack<OpenMPContextFrame> context{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_target,
+ /*isPartial=*/true};
bindEntryBlockArgs(converter, targetOp, args);
if (HostEvalInfo *hostEvalInfo = getHostEvalInfoStackTop(converter))
hostEvalInfo->bindOperands(argIface.getHostEvalBlockArgs());
@@ -5340,6 +5349,9 @@ static mlir::omp::DistributeOp genCompositeDistributeParallelDo(
parallelArgs.reduction.vars = parallelClauseOps.reductionVars;
genParallelOp(converter, symTable, semaCtx, eval, loc, queue, parallelItem,
parallelClauseOps, parallelArgs, &dsp, /*isComposite=*/true);
+ mlir::SaveStateStack<OpenMPContextFrame> context{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_parallel,
+ /*isPartial=*/true};
// Clause processing.
mlir::omp::DistributeOperands distributeClauseOps;
@@ -5410,6 +5422,9 @@ static mlir::omp::DistributeOp genCompositeDistributeParallelDoSimd(
genParallelOp(converter, symTable, semaCtx, eval, loc, queue, parallelItem,
parallelClauseOps, parallelArgs, ¶llelItemDSP,
/*isComposite=*/true);
+ mlir::SaveStateStack<OpenMPContextFrame> context{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_parallel,
+ /*isPartial=*/true};
// Clause processing.
// Use a shared cache so that both wsloop and simd produce the same SSA
@@ -7502,7 +7517,7 @@ static void genMetadirective(lower::AbstractConverter &converter,
fir::FirOpBuilder &builder = converter.getFirOpBuilder();
llvm::SmallVector<llvm::omp::TraitProperty, 8> constructTraits;
- collectEnclosingConstructTraits(converter, eval, constructTraits);
+ collectEnclosingConstructTraits(converter, &eval, constructTraits);
semantics::omp::OmpVariantMatchContext ompCtx =
makeVariantMatchContext(builder.getModule(), constructTraits);
diff --git a/flang/lib/Lower/OpenMP/Utils.cpp b/flang/lib/Lower/OpenMP/Utils.cpp
index 634ced155b0633..8253954d76b631 100644
--- a/flang/lib/Lower/OpenMP/Utils.cpp
+++ b/flang/lib/Lower/OpenMP/Utils.cpp
@@ -1473,7 +1473,7 @@ semantics::omp::OmpVariantMatchContext makeVariantMatchContext(
}
void collectEnclosingConstructTraits(
- AbstractConverter &converter, const pft::Evaluation &evaluation,
+ AbstractConverter &converter, const pft::Evaluation *evaluation,
llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits) {
llvm::SmallVector<const OpenMPContextFrame *, 4> frames;
converter.getStateStack().stackWalk<OpenMPContextFrame>(
@@ -1485,8 +1485,9 @@ void collectEnclosingConstructTraits(
llvm::SmallVector<bool, 4> usedFrames(frames.size(), false);
llvm::SmallVector<const pft::Evaluation *, 8> ancestors;
- for (const pft::Evaluation *parent = evaluation.parentConstruct; parent;
- parent = parent->parentConstruct) {
+ for (const pft::Evaluation *parent = evaluation ? evaluation->parentConstruct
+ : nullptr;
+ parent; parent = parent->parentConstruct) {
ancestors.push_back(parent);
}
std::reverse(ancestors.begin(), ancestors.end());
@@ -1499,25 +1500,37 @@ void collectEnclosingConstructTraits(
if (!omp)
continue;
llvm::omp::Directive directive{parser::omp::GetOmpDirectiveName(*omp).v};
+ // An ancestor supplies the full source context, including constituents
+ // whose bodies also have active frames. Count each construct only once.
+ for (auto [index, frame] : llvm::enumerate(frames))
+ if (&frame->evaluation == ancestor)
+ usedFrames[index] = true;
if (directive != llvm::omp::Directive::OMPD_metadirective) {
append(directive);
continue;
}
- for (auto [index, frame] : llvm::enumerate(frames)) {
- if (&frame->evaluation == ancestor) {
+ for (const OpenMPContextFrame *frame : frames) {
+ if (&frame->evaluation == ancestor && !frame->isPartial) {
append(frame->directive);
- usedFrames[index] = true;
break;
}
}
}
- // Active source contexts may not appear in the ancestor chain. This occurs
- // for a loop-associated metadirective that owns a following sibling and for
- // an atomic expression lowered from the atomic evaluation itself.
- for (auto [index, frame] : llvm::enumerate(frames))
- if (!usedFrames[index])
- append(frame->directive);
+ // Include entered constituents while their own evaluation is current, e.g.
+ // PARALLEL when lowering the bounds of PARALLEL DO. Complete replacement
+ // frames already describe all constituents of a selected metadirective.
+ for (auto [index, frame] : llvm::enumerate(frames)) {
+ if (usedFrames[index])
+ continue;
+ if (frame->isPartial &&
+ llvm::any_of(
+ frames, [frameEval = &frame->evaluation](const auto *other) {
+ return !other->isPartial && &other->evaluation == frameEval;
+ }))
+ continue;
+ append(frame->directive);
+ }
}
const semantics::Symbol *
diff --git a/flang/lib/Lower/OpenMP/Utils.h b/flang/lib/Lower/OpenMP/Utils.h
index 7b89220f432f65..618a357834a9c8 100644
--- a/flang/lib/Lower/OpenMP/Utils.h
+++ b/flang/lib/Lower/OpenMP/Utils.h
@@ -58,11 +58,14 @@ class OpenMPContextFrame
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(OpenMPContextFrame)
OpenMPContextFrame(const pft::Evaluation &evaluation,
- llvm::omp::Directive directive)
- : evaluation{evaluation}, directive{directive} {}
+ llvm::omp::Directive directive, bool isPartial = false)
+ : evaluation{evaluation}, directive{directive}, isPartial{isPartial} {}
const pft::Evaluation &evaluation;
llvm::omp::Directive directive;
+ // A partial frame records an entered constituent of a combined directive.
+ // A complete frame supplies the source context when it is absent from PFT.
+ bool isPartial;
};
struct DeclareTargetCaptureInfo {
@@ -278,9 +281,10 @@ std::optional<llvm::SmallVector<mlir::Value>> getIteratorElementIndices(
/// Collect the source OpenMP constructs enclosing \p evaluation in
/// outermost-to-innermost order. Active metadirective replacements substitute
-/// for their METADIRECTIVE source constructs.
+/// for their METADIRECTIVE source constructs. A null evaluation uses only
+/// active context frames.
void collectEnclosingConstructTraits(
- AbstractConverter &converter, const pft::Evaluation &evaluation,
+ AbstractConverter &converter, const pft::Evaluation *evaluation,
llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits);
/// Return true when \p module is being compiled for an AMDGPU device or all of
diff --git a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90 b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
new file mode 100644
index 00000000000000..30cc4e7d55f4be
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
@@ -0,0 +1,120 @@
+! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir %s -o - | \
+! RUN: FileCheck %s
+
+module loop_context
+contains
+ pure integer function parallel_bound(n)
+ integer, intent(in) :: n
+ parallel_bound = n
+ end function
+
+ pure integer function do_bound(n)
+ integer, intent(in) :: n
+ do_bound = n
+ end function
+
+ pure integer function bound(n)
+ integer, intent(in) :: n
+ !$omp declare variant(parallel_bound) match(construct={parallel})
+ !$omp declare variant(do_bound) match(construct={do})
+ bound = n
+ end function
+
+ ! NUM_THREADS is evaluated outside PARALLEL. The DO bound is inside
+ ! PARALLEL but outside DO, and the loop body is inside both constructs.
+ ! CHECK-LABEL: func.func @_QMloop_contextPcombined(
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: omp.parallel
+ ! CHECK-NOT: fir.call @_QMloop_contextPbound(
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine combined(n, a)
+ integer, intent(in) :: n
+ integer :: a(n), i
+ !$omp parallel do num_threads(bound(n))
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ end subroutine
+
+ ! Explicit nesting must select the same variants as the combined spelling.
+ ! CHECK-LABEL: func.func @_QMloop_contextPexplicit_nest(
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: omp.parallel
+ ! CHECK-NOT: fir.call @_QMloop_contextPbound(
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine explicit_nest(n, a)
+ integer, intent(in) :: n
+ integer :: a(n), i
+ !$omp parallel num_threads(bound(n))
+ !$omp do
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ !$omp end parallel
+ end subroutine
+
+ ! Composite lowering creates PARALLEL separately from the loop wrappers.
+ ! CHECK-LABEL: func.func @_QMloop_contextPcomposite(
+ ! CHECK: omp.teams
+ ! CHECK: omp.parallel
+ ! CHECK-NOT: fir.call @_QMloop_contextPbound(
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.distribute
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine composite(n, a)
+ integer, intent(in) :: n
+ integer :: a(n), i
+ !$omp teams distribute parallel do
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ end subroutine
+
+ ! The SIMD composite follows a separate lowering path.
+ ! CHECK-LABEL: func.func @_QMloop_contextPcomposite_simd(
+ ! CHECK: omp.teams
+ ! CHECK: omp.parallel
+ ! CHECK-NOT: fir.call @_QMloop_contextPbound(
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.distribute
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.simd
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine composite_simd(n, a)
+ integer, intent(in) :: n
+ integer :: a(n), i
+ !$omp teams distribute parallel do simd
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ end subroutine
+
+ ! A standalone DO does not contribute its own context to its bounds.
+ ! CHECK-LABEL: func.func @_QMloop_contextPstandalone_do(
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine standalone_do(n, a)
+ integer, intent(in) :: n
+ integer :: a(n), i
+ !$omp do
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ end subroutine
+end module
diff --git a/flang/test/Lower/OpenMP/declare-variant-spec-expr.f90 b/flang/test/Lower/OpenMP/declare-variant-spec-expr.f90
new file mode 100644
index 00000000000000..704314e68696c9
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-variant-spec-expr.f90
@@ -0,0 +1,56 @@
+! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir %s -o - | \
+! RUN: FileCheck %s --implicit-check-not='fir.call @_QMspec_contextPparallel'
+
+module spec_context
+contains
+ ! Specification expressions are lowered before any current evaluation.
+ ! The caller comes first to avoid reusing a previous procedure's evaluation.
+ ! CHECK-LABEL: func.func @_QMspec_contextPfirst(
+ ! CHECK-NOT: fir.call @_QMspec_contextPbound(
+ ! CHECK: fir.call @_QMspec_contextPreplacement(
+ ! CHECK-NOT: fir.call @_QMspec_contextPreplacement(
+ ! CHECK: fir.call @_QMspec_contextPunmatched(
+ ! CHECK-NOT: fir.call @_QMspec_contextPbound(
+ ! CHECK: return
+ subroutine first(n)
+ integer, intent(in) :: n
+ integer :: a(bound(n)), b(unmatched(n))
+ print *, size(a), size(b)
+ end subroutine
+
+ pure integer function bound(n)
+ integer, intent(in) :: n
+ !$omp declare variant(replacement) match(implementation={vendor(llvm)})
+ bound = n
+ end function
+
+ pure integer function unmatched(n)
+ integer, intent(in) :: n
+ !$omp declare variant(parallel_replacement) match(construct={parallel})
+ unmatched = n
+ end function
+
+ pure integer function replacement(n)
+ integer, intent(in) :: n
+ replacement = n
+ end function
+
+ pure integer function parallel_replacement(n)
+ integer, intent(in) :: n
+ parallel_replacement = n
+ end function
+
+ ! A later procedure's prologue also starts with an empty construct context.
+ ! CHECK-LABEL: func.func @_QMspec_contextPlast(
+ ! CHECK-NOT: fir.call @_QMspec_contextPbound(
+ ! CHECK: fir.call @_QMspec_contextPreplacement(
+ ! CHECK-NOT: fir.call @_QMspec_contextPreplacement(
+ ! CHECK: fir.call @_QMspec_contextPunmatched(
+ ! CHECK-NOT: fir.call @_QMspec_contextPbound(
+ ! CHECK: return
+ subroutine last(n)
+ integer, intent(in) :: n
+ integer :: a(bound(n)), b(unmatched(n))
+ print *, size(a), size(b)
+ end subroutine
+end module
diff --git a/flang/test/Lower/OpenMP/metadirective-target-boundary.f90 b/flang/test/Lower/OpenMP/metadirective-target-boundary.f90
index d7963bf7fe4ece..00adbc0ffdba5a 100644
--- a/flang/test/Lower/OpenMP/metadirective-target-boundary.f90
+++ b/flang/test/Lower/OpenMP/metadirective-target-boundary.f90
@@ -21,22 +21,36 @@ subroutine actual_target(n, a)
!$omp end parallel
end subroutine
-! A TARGET selected by a metadirective creates the same context boundary.
-! CHECK-LABEL: func.func @_QPselected_target(
+! A selected TARGET contributes its own construct trait.
+! CHECK-LABEL: func.func @_QPselected_target_present()
! CHECK: omp.parallel
! CHECK: omp.target
-! CHECK-NOT: omp.simd
+! CHECK-NOT: omp.taskyield
+! CHECK: omp.barrier
+! CHECK-NOT: omp.taskyield
! CHECK: return
-subroutine selected_target(n, a)
- integer :: n, i, a(n)
+subroutine selected_target_present()
!$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& when(construct={target}: barrier) default(taskyield)
+ !$omp end metadirective
+ !$omp end parallel
+end subroutine
+
+! The same selected TARGET hides the enclosing PARALLEL trait.
+! CHECK-LABEL: func.func @_QPselected_target_hides_parallel()
+! CHECK: omp.parallel
+! CHECK: omp.target
+! CHECK-NOT: omp.barrier
+! CHECK: omp.taskyield
+! CHECK-NOT: omp.barrier
+! CHECK: return
+subroutine selected_target_hides_parallel()
+ !$omp parallel
+ !$omp begin metadirective default(target)
+ !$omp metadirective &
+ !$omp& when(construct={parallel}: barrier) default(taskyield)
!$omp end metadirective
!$omp end parallel
end subroutine
More information about the flang-commits
mailing list