[clang] [flang] [llvm] [clang][flang][OpenMP] Fix context selector matching and scoring (PR #224431)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 30 15:01:10 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 01/14] [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 02/14] 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
>From 5757c455f1a86780c5f83033f975c098d25aae26 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Thu, 24 Sep 2026 13:23:15 -0500
Subject: [PATCH 03/14] Fix Flang construct contexts during variant lowering
- Count each enclosing loop construct once, including collapsed nests.
- Match selected directive clauses and bounds in the same contexts as
directly written directives.
- Evaluate composite clause expressions in source construct order.
- Add regressions for scoring and static/runtime variant selection.
---
flang/lib/Lower/OpenMP/OpenMP.cpp | 48 +++---
flang/lib/Lower/OpenMP/Utils.cpp | 15 +-
flang/lib/Lower/OpenMP/Utils.h | 11 +-
.../OpenMP/declare-variant-loop-bounds.f90 | 60 +++++++
.../declare-variant-metadirective-context.f90 | 149 ++++++++++++++++++
.../OpenMP/variant-loop-context-depth.f90 | 118 ++++++++++++++
6 files changed, 367 insertions(+), 34 deletions(-)
create mode 100644 flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90
create mode 100644 flang/test/Lower/OpenMP/variant-loop-context-depth.f90
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index 190349bcb5c061..6d8dedd3e0beb1 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -2017,8 +2017,12 @@ static void createBodyOfOp(mlir::Operation &op, const OpWithBodyGenInfo &info,
return {};
}();
+ // A loop body uses the collapsed DO evaluation, but its context belongs to
+ // the owning directive. This also covers intervening code between loops.
+ auto &contextEval =
+ info.outerCollapseEval ? *info.outerCollapseEval : info.eval;
mlir::SaveStateStack<OpenMPContextFrame> context{
- info.converter.getStateStack(), info.eval, info.dir, /*isPartial=*/true};
+ info.converter.getStateStack(), contextEval, info.dir};
// Mark the earliest insertion point.
mlir::Operation *marker = insertMarker(firOpBuilder);
@@ -2171,8 +2175,7 @@ static void genBodyOfTargetDataOp(
genEntryBlock(firOpBuilder, args.asEntryBlockArgs(), dataOp.getRegion());
mlir::SaveStateStack<OpenMPContextFrame> context{
- converter.getStateStack(), eval, llvm::omp::Directive::OMPD_target_data,
- /*isPartial=*/true};
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_target_data};
bindEntryBlockArgs(converter, dataOp, args);
auto argIface = llvm::cast<mlir::omp::BlockArgOpenMPOpInterface>(*dataOp);
llvm::SmallVector<const semantics::Symbol *> sourceUseDeviceAddrSyms{
@@ -2260,8 +2263,7 @@ 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};
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_target};
bindEntryBlockArgs(converter, targetOp, args);
if (HostEvalInfo *hostEvalInfo = getHostEvalInfoStackTop(converter))
hostEvalInfo->bindOperands(argIface.getHostEvalBlockArgs());
@@ -5330,6 +5332,14 @@ static mlir::omp::DistributeOp genCompositeDistributeParallelDo(
ConstructQueue::const_iterator parallelItem = std::next(distributeItem);
ConstructQueue::const_iterator doItem = std::next(parallelItem);
+ // Clause expressions follow source nesting even though the composite
+ // operations place PARALLEL outside DISTRIBUTE.
+ mlir::omp::DistributeOperands distributeClauseOps;
+ genDistributeClauses(converter, semaCtx, stmtCtx, distributeItem->clauses,
+ loc, distributeClauseOps);
+ mlir::SaveStateStack<OpenMPContextFrame> distributeContext{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_distribute};
+
// Create parent omp.parallel first.
mlir::omp::ParallelOperands parallelClauseOps;
llvm::SmallVector<Object> parallelReductionObjects;
@@ -5350,14 +5360,9 @@ static mlir::omp::DistributeOp genCompositeDistributeParallelDo(
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};
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_parallel};
// Clause processing.
- mlir::omp::DistributeOperands distributeClauseOps;
- genDistributeClauses(converter, semaCtx, stmtCtx, distributeItem->clauses,
- loc, distributeClauseOps);
-
mlir::omp::WsloopOperands wsloopClauseOps;
llvm::SmallVector<Object> wsloopReductionObjects;
genWsloopClauses(converter, semaCtx, stmtCtx, doItem->clauses, loc,
@@ -5401,6 +5406,14 @@ static mlir::omp::DistributeOp genCompositeDistributeParallelDoSimd(
ConstructQueue::const_iterator doItem = std::next(parallelItem);
ConstructQueue::const_iterator simdItem = std::next(doItem);
+ // Evaluate DISTRIBUTE clauses before entering its source context and then
+ // PARALLEL, as in the explicit nesting.
+ mlir::omp::DistributeOperands distributeClauseOps;
+ genDistributeClauses(converter, semaCtx, stmtCtx, distributeItem->clauses,
+ loc, distributeClauseOps);
+ mlir::SaveStateStack<OpenMPContextFrame> distributeContext{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_distribute};
+
// Create parent omp.parallel first.
mlir::omp::ParallelOperands parallelClauseOps;
llvm::SmallVector<Object> parallelReductionObjects;
@@ -5423,18 +5436,13 @@ static mlir::omp::DistributeOp genCompositeDistributeParallelDoSimd(
parallelClauseOps, parallelArgs, ¶llelItemDSP,
/*isComposite=*/true);
mlir::SaveStateStack<OpenMPContextFrame> context{
- converter.getStateStack(), eval, llvm::omp::Directive::OMPD_parallel,
- /*isPartial=*/true};
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_parallel};
// Clause processing.
// Use a shared cache so that both wsloop and simd produce the same SSA
// values for array/box reduction variables. See genCompositeDoSimd.
llvm::DenseMap<const semantics::Symbol *, mlir::Value> reductionVarCache;
- mlir::omp::DistributeOperands distributeClauseOps;
- genDistributeClauses(converter, semaCtx, stmtCtx, distributeItem->clauses,
- loc, distributeClauseOps);
-
mlir::omp::WsloopOperands wsloopClauseOps;
llvm::SmallVector<Object> wsloopReductionObjects;
genWsloopClauses(converter, semaCtx, stmtCtx, doItem->clauses, loc,
@@ -7745,7 +7753,8 @@ static void genMetadirective(lower::AbstractConverter &converter,
TODO(variantLoc, "THREADPRIVATE loop iteration variable in "
"loop-associated METADIRECTIVE variant");
mlir::SaveStateStack<OpenMPContextFrame> context{
- converter.getStateStack(), eval, spec->DirId()};
+ converter.getStateStack(), eval, spec->DirId(),
+ /*isReplacement=*/true};
genOMPDispatch(converter, symTable, semaCtx, eval, variantLoc, queue,
queue.begin(), dsaGuard.getMarkedSymbols());
return;
@@ -7761,7 +7770,8 @@ static void genMetadirective(lower::AbstractConverter &converter,
if (consumesBody) {
mlir::SaveStateStack<OpenMPContextFrame> context{
- converter.getStateStack(), eval, spec->DirId()};
+ converter.getStateStack(), eval, spec->DirId(),
+ /*isReplacement=*/true};
genOMPDispatch(converter, symTable, semaCtx, eval, variantLoc, queue,
queue.begin());
} else {
diff --git a/flang/lib/Lower/OpenMP/Utils.cpp b/flang/lib/Lower/OpenMP/Utils.cpp
index 8253954d76b631..8b489b500e8b51 100644
--- a/flang/lib/Lower/OpenMP/Utils.cpp
+++ b/flang/lib/Lower/OpenMP/Utils.cpp
@@ -1510,7 +1510,7 @@ void collectEnclosingConstructTraits(
continue;
}
for (const OpenMPContextFrame *frame : frames) {
- if (&frame->evaluation == ancestor && !frame->isPartial) {
+ if (&frame->evaluation == ancestor && frame->isReplacement) {
append(frame->directive);
break;
}
@@ -1518,16 +1518,11 @@ void collectEnclosingConstructTraits(
}
// 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.
+ // PARALLEL when lowering the bounds of PARALLEL DO. A selected directive
+ // contributes here only through its entered constituents, so its own clause
+ // expressions have the same context as a directly written directive's.
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;
- }))
+ if (usedFrames[index] || frame->isReplacement)
continue;
append(frame->directive);
}
diff --git a/flang/lib/Lower/OpenMP/Utils.h b/flang/lib/Lower/OpenMP/Utils.h
index 618a357834a9c8..2de5161f760cd2 100644
--- a/flang/lib/Lower/OpenMP/Utils.h
+++ b/flang/lib/Lower/OpenMP/Utils.h
@@ -58,14 +58,15 @@ class OpenMPContextFrame
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(OpenMPContextFrame)
OpenMPContextFrame(const pft::Evaluation &evaluation,
- llvm::omp::Directive directive, bool isPartial = false)
- : evaluation{evaluation}, directive{directive}, isPartial{isPartial} {}
+ llvm::omp::Directive directive, bool isReplacement = false)
+ : evaluation{evaluation}, directive{directive},
+ isReplacement{isReplacement} {}
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;
+ // A replacement records the selected directive for nested evaluations.
+ // Other frames record constructs whose regions have already been entered.
+ bool isReplacement;
};
struct DeclareTargetCaptureInfo {
diff --git a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90 b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
index 30cc4e7d55f4be..c3301c294c6ad3 100644
--- a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
@@ -117,4 +117,64 @@ subroutine standalone_do(n, a)
a(i) = bound(n)
end do
end subroutine
+
+ pure integer function cpu_count(n)
+ integer, intent(in) :: n
+ cpu_count = n
+ end function
+
+ pure integer function scored_count(n)
+ integer, intent(in) :: n
+ scored_count = n + 1
+ end function
+
+ pure integer function thread_count(n)
+ integer, intent(in) :: n
+ !$omp declare variant(cpu_count) match(device={kind(cpu)})
+ !$omp declare variant(scored_count) &
+ !$omp& match(implementation={vendor(score(3): llvm)})
+ thread_count = n
+ end function
+
+ ! DIST_SCHEDULE sees only TEAMS. NUM_THREADS sees TEAMS, DISTRIBUTE, so
+ ! CPU scores 5 and beats the vendor score of 4. The bound also sees PARALLEL.
+ ! CHECK-LABEL: func.func @_QMloop_contextPcomposite_clauses(
+ ! CHECK: omp.teams
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: fir.call @_QMloop_contextPcpu_count(
+ ! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.distribute
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine composite_clauses(n, a)
+ integer :: n, a(n), i
+ !$omp teams distribute parallel do dist_schedule(static, bound(n)) &
+ !$omp& num_threads(thread_count(n))
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ end subroutine
+
+ ! The SIMD composite must evaluate its clauses in the same source contexts.
+ ! CHECK-LABEL: func.func @_QMloop_contextPcomposite_simd_clauses(
+ ! CHECK: omp.teams
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: fir.call @_QMloop_contextPcpu_count(
+ ! CHECK: omp.parallel
+ ! 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_clauses(n, a)
+ integer :: n, a(n), i
+ !$omp teams distribute parallel do simd dist_schedule(static, bound(n)) &
+ !$omp& num_threads(thread_count(n))
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ end subroutine
end module
diff --git a/flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90 b/flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90
new file mode 100644
index 00000000000000..2e5011167c9ee9
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90
@@ -0,0 +1,149 @@
+! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir %s -o - | \
+! RUN: FileCheck %s
+! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-fir %s -o - | \
+! RUN: FileCheck %s
+
+module selected_context
+contains
+ pure integer function in_parallel(n)
+ integer, intent(in) :: n
+ in_parallel = n + 1
+ end function
+
+ pure integer function in_do(n)
+ integer, intent(in) :: n
+ in_do = n + 2
+ end function
+
+ pure integer function value(n)
+ integer, intent(in) :: n
+ !$omp declare variant(in_parallel) match(construct={parallel})
+ !$omp declare variant(in_do) match(construct={do})
+ value = n
+ end function
+
+ ! NUM_THREADS is evaluated outside PARALLEL for either spelling. The body
+ ! sees PARALLEL, and leaving the region restores the enclosing context.
+ ! CHECK-LABEL: func.func @_QMselected_contextPdirect_parallel(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMselected_contextPin_parallel(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine direct_parallel(n, a)
+ integer :: n, a(n)
+ !$omp parallel num_threads(value(n))
+ a(1) = value(n)
+ !$omp end parallel
+ a(n) = value(n)
+ end subroutine
+
+ ! CHECK-LABEL: func.func @_QMselected_contextPstatic_parallel(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMselected_contextPin_parallel(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine static_parallel(n, a)
+ integer :: n, a(n)
+ !$omp begin metadirective &
+ !$omp& when(implementation={vendor(llvm)}: parallel num_threads(value(n))) &
+ !$omp& otherwise(nothing)
+ a(1) = value(n)
+ !$omp end metadirective
+ a(n) = value(n)
+ end subroutine
+
+ ! Each runtime arm uses its own context. NOTHING adds no construct.
+ ! CHECK-LABEL: func.func @_QMselected_contextPruntime_parallel(
+ ! CHECK: fir.if
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMselected_contextPin_parallel(
+ ! CHECK: } else {
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine runtime_parallel(flag, n, a)
+ logical :: flag
+ integer :: n, a(n)
+ !$omp begin metadirective &
+ !$omp& when(user={condition(flag)}: parallel num_threads(value(n))) &
+ !$omp& otherwise(nothing)
+ a(1) = value(n)
+ !$omp end metadirective
+ a(n) = value(n)
+ end subroutine
+
+ ! A loop's bounds exclude its own DO construct, while its body includes it.
+ ! CHECK-LABEL: func.func @_QMselected_contextPdirect_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine direct_do(n, a)
+ integer :: n, a(n), i
+ !$omp do
+ do i = 1, value(n)
+ a(i) = value(n)
+ end do
+ a(n) = value(n)
+ end subroutine
+
+ ! The following DO remains a sibling of the standalone metadirective in
+ ! the source tree, but its body must still see the selected construct.
+ ! CHECK-LABEL: func.func @_QMselected_contextPstatic_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine static_do(n, a)
+ integer :: n, a(n), i
+ !$omp metadirective when(implementation={vendor(llvm)}: do) &
+ !$omp& otherwise(nothing)
+ do i = 1, value(n)
+ a(i) = value(n)
+ end do
+ a(n) = value(n)
+ end subroutine
+
+ ! CHECK-LABEL: func.func @_QMselected_contextPruntime_do(
+ ! CHECK: fir.if
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: } else {
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine runtime_do(flag, n, a)
+ logical :: flag
+ integer :: n, a(n), i
+ !$omp metadirective when(user={condition(flag)}: do) otherwise(nothing)
+ do i = 1, value(n)
+ a(i) = value(n)
+ end do
+ a(n) = value(n)
+ end subroutine
+
+ ! The delimited form owns the loop in the source tree.
+ ! CHECK-LABEL: func.func @_QMselected_contextPdelimited_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine delimited_do(n, a)
+ integer :: n, a(n), i
+ !$omp begin metadirective &
+ !$omp& when(implementation={vendor(llvm)}: do) otherwise(nothing)
+ do i = 1, value(n)
+ a(i) = value(n)
+ end do
+ !$omp end metadirective
+ a(n) = value(n)
+ end subroutine
+end module
diff --git a/flang/test/Lower/OpenMP/variant-loop-context-depth.f90 b/flang/test/Lower/OpenMP/variant-loop-context-depth.f90
new file mode 100644
index 00000000000000..20a63d47f090db
--- /dev/null
+++ b/flang/test/Lower/OpenMP/variant-loop-context-depth.f90
@@ -0,0 +1,118 @@
+! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir %s -o - | \
+! RUN: FileCheck %s --implicit-check-not=omp.taskyield \
+! RUN: --implicit-check-not='fir.call @_QMloop_depthPdo_value'
+! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-fir %s -o - | \
+! RUN: FileCheck %s --implicit-check-not=omp.taskyield \
+! RUN: --implicit-check-not='fir.call @_QMloop_depthPdo_value'
+
+module loop_depth
+contains
+ integer function do_value()
+ do_value = 1
+ end function
+
+ integer function scored_value()
+ scored_value = 2
+ end function
+
+ integer function value()
+ !$omp declare variant(do_value) match(construct={do})
+ !$omp declare variant(scored_value) &
+ !$omp& match(implementation={vendor(score(3): llvm)})
+ value = 0
+ end function
+
+ ! At depth one CPU scores 3, below the user selector's score of 4.
+ ! Counting the DO twice would select TASKYIELD instead.
+ ! CHECK-LABEL: func.func @_QMloop_depthPbare_do(
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.taskwait
+ ! CHECK: return
+ subroutine bare_do(n)
+ integer :: n, i
+ !$omp do
+ do i = 1, n
+ !$omp metadirective when(device={kind(cpu)}: taskyield) &
+ !$omp& when(user={condition(score(3): .true.)}: taskwait)
+ end do
+ end subroutine
+
+ ! DO at position two scores 3, below the vendor selector's score of 4.
+ ! The same context must be used for both kinds of variant selection.
+ ! CHECK-LABEL: func.func @_QMloop_depthPcombined(
+ ! CHECK: omp.parallel
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.taskwait
+ ! CHECK: fir.call @_QMloop_depthPscored_value()
+ ! CHECK: return
+ subroutine combined(n, a)
+ integer :: n, i, a(n)
+ !$omp parallel do
+ do i = 1, n
+ !$omp metadirective when(construct={do}: taskyield) &
+ !$omp& when(implementation={vendor(score(3): llvm)}: taskwait)
+ a(i) = value()
+ end do
+ end subroutine
+
+ ! Intervening code is outside the innermost DO evaluation, but shares the
+ ! same OpenMP construct context as the collapsed loop body.
+ ! CHECK-LABEL: func.func @_QMloop_depthPcollapsed(
+ ! CHECK: omp.parallel
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_depthPscored_value()
+ ! CHECK: fir.call @_QMloop_depthPscored_value()
+ ! CHECK: fir.call @_QMloop_depthPscored_value()
+ ! CHECK: return
+ subroutine collapsed(n, a)
+ integer :: n, i, j, a(n, n)
+ !$omp parallel do collapse(2)
+ do i = 1, n
+ a(i, 1) = value()
+ do j = 1, n
+ a(i, j) = value()
+ end do
+ a(i, n) = value()
+ end do
+ end subroutine
+
+ ! A nested TASK contributes a position, but must not move the enclosing DO.
+ ! CHECK-LABEL: func.func @_QMloop_depthPnested_task(
+ ! CHECK: omp.parallel
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.task
+ ! CHECK: omp.taskwait
+ ! CHECK: return
+ subroutine nested_task(n)
+ integer :: n, i
+ !$omp parallel do
+ do i = 1, n
+ !$omp task
+ !$omp metadirective when(construct={do}: taskyield) &
+ !$omp& when(implementation={vendor(score(3): llvm)}: taskwait)
+ !$omp end task
+ end do
+ end subroutine
+
+ ! TARGET hides the outer PARALLEL. TARGET, PARALLEL, DO give depth three:
+ ! CPU scores 9 and loses to the user selector's score of 10.
+ ! CHECK-LABEL: func.func @_QMloop_depthPtarget_depth(
+ ! CHECK: omp.parallel
+ ! CHECK: omp.target
+ ! CHECK: omp.parallel
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.taskwait
+ ! CHECK: return
+ subroutine target_depth(n)
+ integer :: n, i
+ !$omp parallel
+ !$omp target
+ !$omp parallel do
+ do i = 1, n
+ !$omp metadirective when(device={kind(cpu)}: taskyield) &
+ !$omp& when(user={condition(score(9): .true.)}: taskwait)
+ end do
+ !$omp end target
+ !$omp end parallel
+ end subroutine
+end module
>From f356de2d0eed7fbcd72612fd2a201985e9253a1c Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Thu, 24 Sep 2026 13:43:41 -0500
Subject: [PATCH 04/14] Preserve declaration identity in user conditions
Use canonical AST profiles instead of printed expressions for condition
subset checks. Preserve original expressions through template substitution
and PCH serialization so distinct declarations remain distinguishable.
Add regression coverage for template parameters, scoped declarations,
and PCH round trips.
---
clang/include/clang/AST/OpenMPClause.h | 4 +-
clang/lib/AST/OpenMPClause.cpp | 21 +++-
clang/lib/Parse/ParseOpenMP.cpp | 12 +-
clang/lib/Serialization/ASTReader.cpp | 2 +
clang/lib/Serialization/ASTWriter.cpp | 3 +
...declare_variant_condition_declarations.cpp | 107 ++++++++++++++++++
.../declare_variant_condition_identity_pch.c | 6 +-
7 files changed, 138 insertions(+), 17 deletions(-)
create mode 100644 clang/test/OpenMP/declare_variant_condition_declarations.cpp
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 18f1b404445b8c..dd48ad6f375925 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -10057,12 +10057,14 @@ 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), user-condition identity, and extensions.
+ /// (which accepts anything) and extensions.
StringRef RawString;
};
struct OMPTraitSelector {
Expr *ScoreOrCondition = nullptr;
+ /// Preserve the condition before template substitution for subset checks.
+ Expr *OriginalCondition = nullptr;
llvm::omp::TraitSelector Kind = llvm::omp::TraitSelector::invalid;
SmallVector<OMPTraitProperty, 1> Properties;
};
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 077b81c47cae49..3a5d547b944871 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -20,6 +20,7 @@
#include "clang/Basic/LLVM.h"
#include "clang/Basic/OpenMPKinds.h"
#include "clang/Basic/TargetInfo.h"
+#include "llvm/ADT/FoldingSet.h"
#include "llvm/ADT/Sequence.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/Support/ErrorHandling.h"
@@ -3126,14 +3127,28 @@ void OMPTraitInfo::getAsVariantMatchInfo(ASTContext &ASTCtx,
TraitProperty::user_condition_unknown &&
"Ill-formed user condition, expected unknown trait property!");
+ // Profile the original expression so equally valued template arguments
+ // do not erase distinctions between different parameters. Canonical
+ // profiles also distinguish declarations with the same printed name.
+ // Compute this after loading the AST: profiles contain local pointers
+ // and cannot be compared with profiles serialized by another process.
+ assert(Selector.OriginalCondition && "missing original condition");
+ llvm::FoldingSetNodeID ID;
+ Selector.OriginalCondition->IgnoreParenImpCasts()->Profile(
+ ID, ASTCtx, /*Canonical=*/true);
+ llvm::FoldingSetNodeIDRef Profile = ID.Intern(ASTCtx.getAllocator());
+ // Keep the full profile to avoid hash collisions in subset checks.
+ StringRef ConditionIdentity(
+ reinterpret_cast<const char *>(Profile.data()),
+ Profile.size() * sizeof(unsigned));
+
if (std::optional<APSInt> CondVal =
Selector.ScoreOrCondition->getIntegerConstantExpr(ASTCtx))
VMI.addTrait(CondVal->isZero() ? TraitProperty::user_condition_false
: TraitProperty::user_condition_true,
- Selector.Properties.front().RawString);
+ ConditionIdentity);
else
- VMI.addTrait(TraitProperty::user_condition_false,
- Selector.Properties.front().RawString);
+ VMI.addTrait(TraitProperty::user_condition_false, ConditionIdentity);
continue;
}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index c6a358f8af823e..8b93aca37b6b63 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -28,7 +28,6 @@
#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>
@@ -1151,16 +1150,9 @@ void Parser::parseOMPContextSelector(
if (!Condition.isUsable())
return FinishSelector();
TISelector.ScoreOrCondition = Condition.get();
- // 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.OriginalCondition = Condition.get();
TISelector.Properties.push_back(
- {TraitProperty::user_condition_unknown,
- StringRef(ConditionText).copy(Context.getAllocator())});
+ {TraitProperty::user_condition_unknown, "<condition>"});
return;
}
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 8910b28e6c1659..ef1c8320bcf5cc 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -13134,6 +13134,8 @@ OMPTraitInfo *ASTRecordReader::readOMPTraitInfo() {
Selector.ScoreOrCondition = nullptr;
if (readBool())
Selector.ScoreOrCondition = readExprRef();
+ if (readBool())
+ Selector.OriginalCondition = readExprRef();
Selector.Properties.resize(readUInt32());
for (auto &Property : Selector.Properties) {
Property.Kind = readEnum<llvm::omp::TraitProperty>();
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index a2db5b1a9bd1e6..4e14f27c28cb04 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8977,6 +8977,9 @@ void ASTRecordWriter::writeOMPTraitInfo(const OMPTraitInfo *TI) {
writeBool(Selector.ScoreOrCondition);
if (Selector.ScoreOrCondition)
writeExprRef(Selector.ScoreOrCondition);
+ writeBool(Selector.OriginalCondition);
+ if (Selector.OriginalCondition)
+ writeExprRef(Selector.OriginalCondition);
writeUInt32(Selector.Properties.size());
for (const auto &Property : Selector.Properties) {
writeEnum(Property.Kind);
diff --git a/clang/test/OpenMP/declare_variant_condition_declarations.cpp b/clang/test/OpenMP/declare_variant_condition_declarations.cpp
new file mode 100644
index 00000000000000..cbdbab8509e2b8
--- /dev/null
+++ b/clang/test/OpenMP/declare_variant_condition_declarations.cpp
@@ -0,0 +1,107 @@
+// RUN: split-file %s %t
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -include %t/conditions.h \
+// RUN: %t/use.cpp -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.pch
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -include-pch %t/conditions.pch \
+// RUN: %t/use.cpp -emit-llvm -o - | FileCheck %s
+
+// Conditions naming different parameters remain distinct even when their
+// arguments have the same value. Exercise an attribute instantiated before
+// serialization as well as attributes instantiated after loading the PCH.
+// CHECK-LABEL: define{{.*}} i32 @instantiated_before_pch()
+// CHECK: call i32 @high()
+// CHECK: ret i32
+// CHECK-LABEL: define{{.*}} i32 @distinct_parameters()
+// CHECK: call i32 @high()
+// CHECK: call i32 @high()
+// CHECK: ret i32
+
+// Parameter names do not affect identity when their positions agree.
+// CHECK-LABEL: define{{.*}} i32 @renamed_parameters()
+// CHECK: call i32 @low()
+// CHECK: ret i32
+
+// Equal-valued enumerators in different scopes are different declarations.
+// CHECK-LABEL: define{{.*}} i32 @different_scopes()
+// CHECK: call i32 @high()
+// CHECK: ret i32
+
+// Compare a condition loaded from a PCH with one parsed in this process.
+// They name the same declaration and must still form a subset relationship.
+// CHECK-LABEL: define{{.*}} i32 @same_declaration()
+// CHECK: call i32 @low()
+// CHECK: ret i32
+
+//--- conditions.h
+extern "C" int high();
+extern "C" int low();
+
+#pragma omp declare variant(high) \
+ match(user = {condition(N)}, \
+ implementation = {vendor(score(100) : llvm)})
+template <int N, int M> int distinct();
+#pragma omp declare variant(low) \
+ match(user = {condition(N)}, device = {kind(cpu)}, \
+ implementation = {vendor(score(1) : llvm)})
+template <int M, int N> int distinct();
+
+extern "C" int instantiated_before_pch() { return distinct<1, 1>(); }
+
+#pragma omp declare variant(high) \
+ match(user = {condition(N)}, \
+ implementation = {vendor(score(100) : llvm)})
+template <int N> int renamed();
+#pragma omp declare variant(low) \
+ match(user = {condition(A)}, device = {kind(cpu)}, \
+ implementation = {vendor(score(1) : llvm)})
+template <int A> int renamed();
+
+namespace First {
+enum { A = 1 };
+extern "C" {
+#pragma omp declare variant(high) \
+ match(user = {condition(A)}, \
+ implementation = {vendor(score(100) : llvm)})
+int scoped_base();
+}
+}
+
+enum { Shared = 1 };
+#pragma omp declare variant(high) \
+ match(user = {condition(Shared)}, \
+ implementation = {vendor(score(100) : llvm)})
+int shared_base();
+
+//--- use.cpp
+namespace Second {
+enum { A = 1 };
+extern "C" {
+#pragma omp declare variant(low) \
+ match(user = {condition(A)}, device = {kind(cpu)}, \
+ implementation = {vendor(score(1) : llvm)})
+int scoped_base();
+}
+}
+
+#pragma omp declare variant(low) \
+ match(user = {condition(Shared)}, device = {kind(cpu)}, \
+ implementation = {vendor(score(1) : llvm)})
+int shared_base();
+
+extern "C" int distinct_parameters() {
+ return distinct<1, 2>() + distinct<1, 1>();
+}
+
+extern "C" int renamed_parameters() { return renamed<1>(); }
+
+extern "C" int different_scopes() {
+ return Second::scoped_base();
+}
+
+extern "C" int same_declaration() {
+ return shared_base();
+}
diff --git a/clang/test/OpenMP/declare_variant_condition_identity_pch.c b/clang/test/OpenMP/declare_variant_condition_identity_pch.c
index 166612af3a9335..57d2c16ba8bf21 100644
--- a/clang/test/OpenMP/declare_variant_condition_identity_pch.c
+++ b/clang/test/OpenMP/declare_variant_condition_identity_pch.c
@@ -12,9 +12,9 @@
// 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 expanded identity of folded user
-// conditions. Distinct conditions do not form a subset relationship, while
-// identical conditions do.
+// Verify that serialization preserves the expressions of folded user
+// conditions. Expanded expressions identify distinct conditions, while
+// ignoring outer parentheses preserves identical conditions.
// CHECK-LABEL: define{{.*}} void @test_conditions()
// CHECK: call void @condition_high_variant()
>From 9dedfee92fb852d7f22ff866a3fdd03db740fd4f Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Fri, 25 Sep 2026 01:33:11 -0500
Subject: [PATCH 05/14] Fix construct contexts for SIMD clause expressions
Enter DO while lowering SIMD clauses in DO SIMD and DISTRIBUTE PARALLEL
DO SIMD, then restore the context before lowering loop bounds.
Check SIMD clause selection, TEAMS in DIST_SCHEDULE, and the absence of
PARALLEL in NUM_THREADS operands.
---
flang/lib/Lower/OpenMP/OpenMP.cpp | 20 ++++-
.../OpenMP/declare-variant-loop-bounds.f90 | 74 +++++++++++++++++--
2 files changed, 85 insertions(+), 9 deletions(-)
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index 6d8dedd3e0beb1..5e568c1e2965a6 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -5450,8 +5450,14 @@ static mlir::omp::DistributeOp genCompositeDistributeParallelDoSimd(
mlir::omp::SimdOperands simdClauseOps;
llvm::SmallVector<Object> simdReductionObjects;
- genSimdClauses(converter, semaCtx, simdItem->clauses, loc, simdClauseOps,
- simdReductionObjects, &reductionVarCache);
+ {
+ // SIMD clause expressions are inside DO, but loop control remains outside
+ // the loop constructs.
+ mlir::SaveStateStack<OpenMPContextFrame> doContext{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_do};
+ genSimdClauses(converter, semaCtx, simdItem->clauses, loc, simdClauseOps,
+ simdReductionObjects, &reductionVarCache);
+ }
// Same as genCompositeDoSimd.
if (!simdClauseOps.linearVars.empty()) {
@@ -5598,8 +5604,14 @@ static mlir::omp::WsloopOp genCompositeDoSimd(
mlir::omp::SimdOperands simdClauseOps;
llvm::SmallVector<Object> simdReductionObjects;
- genSimdClauses(converter, semaCtx, simdItem->clauses, loc, simdClauseOps,
- simdReductionObjects, &reductionVarCache);
+ {
+ // SIMD clause expressions are inside DO, but loop control remains outside
+ // the loop constructs.
+ mlir::SaveStateStack<OpenMPContextFrame> doContext{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_do};
+ genSimdClauses(converter, semaCtx, simdItem->clauses, loc, simdClauseOps,
+ simdReductionObjects, &reductionVarCache);
+ }
// omp.simd writes back linear vars unconditionally, causing a race when
// inside a parallel region. Move them to wsloop which has proper last-iter
diff --git a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90 b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
index c3301c294c6ad3..85a301bdfe3a8e 100644
--- a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
@@ -1,8 +1,15 @@
! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir %s -o - | \
-! RUN: FileCheck %s
+! RUN: FileCheck %s --implicit-check-not='fir.call @_QMloop_contextPparallel_count'
+! RUN: %flang_fc1 -fopenmp -fopenmp-version=52 -emit-fir %s -o - | \
+! RUN: FileCheck %s --implicit-check-not='fir.call @_QMloop_contextPparallel_count'
module loop_context
contains
+ pure integer function teams_bound(n)
+ integer, intent(in) :: n
+ teams_bound = n
+ end function
+
pure integer function parallel_bound(n)
integer, intent(in) :: n
parallel_bound = n
@@ -15,6 +22,7 @@ pure integer function do_bound(n)
pure integer function bound(n)
integer, intent(in) :: n
+ !$omp declare variant(teams_bound) match(construct={teams})
!$omp declare variant(parallel_bound) match(construct={parallel})
!$omp declare variant(do_bound) match(construct={do})
bound = n
@@ -128,19 +136,38 @@ pure integer function scored_count(n)
scored_count = n + 1
end function
+ pure integer function parallel_count(n)
+ integer, intent(in) :: n
+ parallel_count = n + 2
+ end function
+
+ pure logical function do_pred(n)
+ integer, intent(in) :: n
+ do_pred = n > 0
+ end function
+
+ pure logical function pred(n)
+ integer, intent(in) :: n
+ !$omp declare variant(do_pred) match(construct={do})
+ pred = n > 0
+ end function
+
pure integer function thread_count(n)
integer, intent(in) :: n
!$omp declare variant(cpu_count) match(device={kind(cpu)})
+ !$omp declare variant(parallel_count) &
+ !$omp& match(construct={parallel}, device={kind(cpu)})
!$omp declare variant(scored_count) &
!$omp& match(implementation={vendor(score(3): llvm)})
thread_count = n
end function
! DIST_SCHEDULE sees only TEAMS. NUM_THREADS sees TEAMS, DISTRIBUTE, so
- ! CPU scores 5 and beats the vendor score of 4. The bound also sees PARALLEL.
+ ! CPU scores 5 and beats the vendor score of 4. A leaked PARALLEL would
+ ! make parallel_count a strict superset of cpu_count. The bound sees PARALLEL.
! CHECK-LABEL: func.func @_QMloop_contextPcomposite_clauses(
! CHECK: omp.teams
- ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: fir.call @_QMloop_contextPteams_bound(
! CHECK: fir.call @_QMloop_contextPcpu_count(
! CHECK: omp.parallel
! CHECK: fir.call @_QMloop_contextPparallel_bound(
@@ -160,9 +187,10 @@ subroutine composite_clauses(n, a)
! The SIMD composite must evaluate its clauses in the same source contexts.
! CHECK-LABEL: func.func @_QMloop_contextPcomposite_simd_clauses(
! CHECK: omp.teams
- ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: fir.call @_QMloop_contextPteams_bound(
! CHECK: fir.call @_QMloop_contextPcpu_count(
! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMloop_contextPdo_pred(
! CHECK: fir.call @_QMloop_contextPparallel_bound(
! CHECK: omp.distribute
! CHECK: omp.wsloop
@@ -172,9 +200,45 @@ subroutine composite_clauses(n, a)
subroutine composite_simd_clauses(n, a)
integer :: n, a(n), i
!$omp teams distribute parallel do simd dist_schedule(static, bound(n)) &
- !$omp& num_threads(thread_count(n))
+ !$omp& num_threads(thread_count(n)) if(simd: pred(n))
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ end subroutine
+
+ ! SIMD clauses see DO, but the following loop bound and post-loop call do not.
+ ! CHECK-LABEL: func.func @_QMloop_contextPdo_simd_clauses(
+ ! CHECK: fir.call @_QMloop_contextPdo_pred(
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.simd
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPpred(
+ ! CHECK: return
+ subroutine do_simd_clauses(n, a, flag)
+ integer :: n, a(n), i
+ logical :: flag
+ !$omp do simd if(simd: pred(n))
do i = 1, bound(n)
a(i) = bound(n)
end do
+ flag = pred(n)
+ end subroutine
+
+ ! The explicit DO/SIMD nesting also exposes DO to the SIMD clause.
+ ! CHECK-LABEL: func.func @_QMloop_contextPnested_simd_clauses(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_contextPdo_pred(
+ ! CHECK: omp.simd
+ ! CHECK: return
+ subroutine nested_simd_clauses(n, a)
+ integer :: n, a(n, n), i, j
+ !$omp do
+ do i = 1, n
+ !$omp simd if(simd: pred(n))
+ do j = 1, n
+ a(i, j) = i + j
+ end do
+ end do
end subroutine
end module
>From 1b3d763e326705a4494ac897b98440f3fd9bba95 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Fri, 25 Sep 2026 01:33:11 -0500
Subject: [PATCH 06/14] Strengthen variant context and ISA serialization tests
Add DO-only calls at each collapsed-loop position so missing and duplicate
DO contexts are both detected. Resolve ISA variant calls after loading a
PCH, with and without the requested target feature.
---
clang/test/OpenMP/declare_variant_isa_pch.c | 30 +++++++++++++++++++
.../OpenMP/variant-loop-context-depth.f90 | 17 ++++++++---
2 files changed, 43 insertions(+), 4 deletions(-)
create mode 100644 clang/test/OpenMP/declare_variant_isa_pch.c
diff --git a/clang/test/OpenMP/declare_variant_isa_pch.c b/clang/test/OpenMP/declare_variant_isa_pch.c
new file mode 100644
index 00000000000000..36364e96b4ef88
--- /dev/null
+++ b/clang/test/OpenMP/declare_variant_isa_pch.c
@@ -0,0 +1,30 @@
+// RUN: split-file %s %t
+// RUN: %clang_cc1 -fopenmp -triple x86_64-unknown-linux -target-feature +avx512f \
+// RUN: -x c-header %t/variants.h -emit-pch -o %t/variants.pch
+// RUN: %clang_cc1 -fopenmp -triple x86_64-unknown-linux -target-feature +avx512f \
+// RUN: -include-pch %t/variants.pch %t/use.c -emit-llvm -o - -verify | \
+// RUN: FileCheck %s --check-prefix=AVX512
+// RUN: %clang_cc1 -fopenmp -triple x86_64-unknown-linux \
+// RUN: -x c-header %t/variants.h -emit-pch -o %t/generic.pch
+// RUN: %clang_cc1 -fopenmp -triple x86_64-unknown-linux \
+// RUN: -include-pch %t/generic.pch %t/use.c -emit-llvm -o - -verify | \
+// RUN: FileCheck %s --check-prefix=GENERIC
+
+// Resolve the call after loading the selector, not while building the PCH.
+// AVX512-LABEL: define{{.*}} void @caller(
+// AVX512: call void @avx_variant()
+// AVX512-NOT: call void @base_isa()
+// AVX512: ret void
+// GENERIC-LABEL: define{{.*}} void @caller(
+// GENERIC: call void @base_isa()
+// GENERIC-NOT: call void @avx_variant()
+// GENERIC: ret void
+
+//--- variants.h
+void avx_variant(void);
+#pragma omp declare variant(avx_variant) match(device={isa(avx512f)})
+void base_isa(void);
+
+//--- use.c
+// expected-no-diagnostics
+void caller(void) { base_isa(); }
diff --git a/flang/test/Lower/OpenMP/variant-loop-context-depth.f90 b/flang/test/Lower/OpenMP/variant-loop-context-depth.f90
index 20a63d47f090db..87ca89d3b70f37 100644
--- a/flang/test/Lower/OpenMP/variant-loop-context-depth.f90
+++ b/flang/test/Lower/OpenMP/variant-loop-context-depth.f90
@@ -22,6 +22,11 @@ integer function value()
value = 0
end function
+ integer function requiring_do()
+ !$omp declare variant(do_value) match(construct={do})
+ requiring_do = 0
+ end function
+
! At depth one CPU scores 3, below the user selector's score of 4.
! Counting the DO twice would select TASKYIELD instead.
! CHECK-LABEL: func.func @_QMloop_depthPbare_do(
@@ -56,23 +61,27 @@ subroutine combined(n, a)
end subroutine
! Intervening code is outside the innermost DO evaluation, but shares the
- ! same OpenMP construct context as the collapsed loop body.
+ ! same OpenMP construct context as the collapsed loop body. The competing
+ ! call rejects a duplicated DO; the DO-only call rejects a missing DO.
! CHECK-LABEL: func.func @_QMloop_depthPcollapsed(
! CHECK: omp.parallel
! CHECK: omp.wsloop
! CHECK: fir.call @_QMloop_depthPscored_value()
+ ! CHECK: fir.call @_QMloop_depthPdo_value()
! CHECK: fir.call @_QMloop_depthPscored_value()
+ ! CHECK: fir.call @_QMloop_depthPdo_value()
! CHECK: fir.call @_QMloop_depthPscored_value()
+ ! CHECK: fir.call @_QMloop_depthPdo_value()
! CHECK: return
subroutine collapsed(n, a)
integer :: n, i, j, a(n, n)
!$omp parallel do collapse(2)
do i = 1, n
- a(i, 1) = value()
+ a(i, 1) = value() + requiring_do()
do j = 1, n
- a(i, j) = value()
+ a(i, j) = value() + requiring_do()
end do
- a(i, n) = value()
+ a(i, n) = value() + requiring_do()
end do
end subroutine
>From 1aa333aaf4a1e91c33e979a5ed5d781f796d6ab4 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Fri, 25 Sep 2026 01:33:12 -0500
Subject: [PATCH 07/14] Normalize Flang user condition identity
Ignore outer parentheses and resolve USE aliases to their ultimate
symbols, preserving original expressions and distinct declarations.
Carry the same identity through static and runtime candidate ranking.
Add coverage for aliases, equal-valued declarations, imported symbols,
and false-condition ranking.
---
flang/lib/Semantics/openmp-utils.cpp | 89 +++++++-
.../declare-variant-condition-identity.f90 | 198 ++++++++++++++++++
2 files changed, 280 insertions(+), 7 deletions(-)
create mode 100644 flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
diff --git a/flang/lib/Semantics/openmp-utils.cpp b/flang/lib/Semantics/openmp-utils.cpp
index 6e311cd2e38f4f..a65b2235445bd5 100644
--- a/flang/lib/Semantics/openmp-utils.cpp
+++ b/flang/lib/Semantics/openmp-utils.cpp
@@ -28,6 +28,7 @@
#include "flang/Evaluate/type.h"
#include "flang/Evaluate/variable.h"
#include "flang/Parser/openmp-utils.h"
+#include "flang/Parser/parse-tree-visitor.h"
#include "flang/Parser/parse-tree.h"
#include "flang/Semantics/expression.h"
#include "flang/Semantics/openmp-directive-sets.h"
@@ -38,6 +39,7 @@
#include "llvm/ADT/APInt.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/DenseSet.h"
+#include "llvm/ADT/FoldingSet.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/SmallVector.h"
@@ -2482,6 +2484,80 @@ void AppendDirectiveContextTraits(llvm::omp::Directive directive,
}
}
+namespace {
+struct ConditionNames {
+ template <typename A> bool Pre(const A &) { return true; }
+ template <typename A> void Post(const A &) {}
+ bool Pre(const parser::Name &name) {
+ if (name.symbol)
+ names.push_back(&name);
+ return false;
+ }
+ llvm::SmallVector<const parser::Name *> names;
+};
+
+// Profile the original expression, not its folded value. In particular, two
+// named constants with the same value still denote different declarations.
+// Retain the cooked expression spelling, except for outer parentheses, blanks
+// outside character literals, and names resolved to their ultimate symbols.
+llvm::StringRef GetConditionIdentity(
+ const parser::ScalarExpr &condition, SemanticsContext &context) {
+ const auto *expr{parser::Unwrap<parser::Expr>(condition)};
+ CHECK(expr);
+ while (
+ const auto *parentheses{std::get_if<parser::Expr::Parentheses>(&expr->u)})
+ expr = &parentheses->v.value();
+
+ ConditionNames visitor;
+ parser::Walk(*expr, visitor);
+ llvm::sort(visitor.names, [](const auto *a, const auto *b) {
+ return a->source.begin() < b->source.begin();
+ });
+
+ llvm::FoldingSetNodeID id;
+ auto addText = [&](parser::CharBlock source) {
+ std::string text;
+ char quote{0};
+ for (const char *p{source.begin()}; p != source.end(); ++p) {
+ char ch{*p};
+ if (quote) {
+ text += ch;
+ if (ch == quote)
+ quote = 0;
+ else if (ch == '\\' &&
+ context.IsEnabled(common::LanguageFeature::BackslashEscapes) &&
+ p + 1 != source.end())
+ text += *++p;
+ } else if (ch == '\'' || ch == '"') {
+ quote = ch;
+ text += ch;
+ } else if (!llvm::isSpace(ch)) {
+ text += ch;
+ }
+ }
+ id.AddString(text);
+ };
+
+ const char *next{expr->source.begin()};
+ for (const parser::Name *name : visitor.names) {
+ if (!expr->source.Contains(name->source) || name->source.begin() < next)
+ continue;
+ addText({next, name->source.begin()});
+ id.AddPointer(&name->symbol->GetUltimate());
+ next = name->source.end();
+ }
+ addText({next, expr->source.end()});
+
+ // Intern the full profile in context-owned storage. Recompute it after
+ // reading a module; symbol pointers are local to this compilation.
+ llvm::FoldingSetNodeIDRef profile{id.getRef()};
+ SourceName saved{context.SaveTempName(
+ std::string{reinterpret_cast<const char *>(profile.data()),
+ profile.size() * sizeof(unsigned)})};
+ return {saved.begin(), saved.size()};
+}
+} // namespace
+
static void AddTraitPropertiesFromSelector(llvm::omp::TraitSet set,
const parser::OmpTraitSelector &selector, llvm::omp::VariantMatchInfo &vmi,
SemanticsContext &semaCtx,
@@ -2508,8 +2584,8 @@ static void AddTraitPropertiesFromSelector(llvm::omp::TraitSet set,
if (!scalarExpr) {
continue;
}
+ llvm::StringRef condition{GetConditionIdentity(*scalarExpr, semaCtx)};
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,
@@ -2520,7 +2596,7 @@ static void AddTraitPropertiesFromSelector(llvm::omp::TraitSet set,
dynamicCond = DynamicUserCondition{scalarExpr, prop.source};
}
vmi.addTrait(set, llvm::omp::TraitProperty::user_condition_unknown,
- llvm::StringRef{prop.source.begin(), prop.source.size()}, scorePtr);
+ condition, scorePtr);
}
return;
}
@@ -2640,8 +2716,7 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet(
staticVMI.UserCondition = {};
llvm::APInt *conditionScorePtr{
conditionScore ? &*conditionScore : nullptr};
- llvm::StringRef conditionSource{
- dynamicCondition->source.begin(), dynamicCondition->source.size()};
+ llvm::StringRef conditionIdentity{rawVMI.UserCondition};
bool hasMatchAny{rawVMI.RequiredTraits.test(unsigned(matchAnyTrait))};
bool hasMatchNone{rawVMI.RequiredTraits.test(unsigned(matchNoneTrait))};
@@ -2656,7 +2731,7 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet(
llvm::omp::VariantMatchInfo conditionTrueVMI{staticVMI};
conditionTrueVMI.addTrait(
- llvm::omp::TraitProperty::user_condition_true, conditionSource,
+ llvm::omp::TraitProperty::user_condition_true, conditionIdentity,
conditionScorePtr);
if (!llvm::omp::isVariantApplicableInContext(
conditionTrueVMI, matchContext)) {
@@ -2669,7 +2744,7 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet(
rankingVMI.addTrait(hasMatchNone
? llvm::omp::TraitProperty::user_condition_false
: llvm::omp::TraitProperty::user_condition_true,
- conditionSource, conditionScorePtr);
+ conditionIdentity, conditionScorePtr);
};
if (hasMatchAny && isStaticVMIApplicable) {
@@ -2680,7 +2755,7 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet(
addConditionTraitForRanking(conditionTrueVMI);
llvm::omp::VariantMatchInfo conditionFalseVMI{staticVMI};
conditionFalseVMI.addTrait(
- llvm::omp::TraitProperty::user_condition_false, conditionSource,
+ llvm::omp::TraitProperty::user_condition_false, conditionIdentity,
conditionScorePtr);
result.candidates.push_back({spec, std::move(conditionTrueVMI),
isExplicit, dynamicCondition});
diff --git a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90 b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
new file mode 100644
index 00000000000000..93c77029df1338
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
@@ -0,0 +1,198 @@
+! RUN: split-file %s %t
+! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -fsyntax-only settings.f90
+! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir variants.f90 -o - | FileCheck %s --check-prefix=LOCAL
+! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir use.f90 -o - | FileCheck %s --check-prefix=IMPORTED
+! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -emit-fir use.f90 -o - | FileCheck %s --check-prefix=IMPORTED
+
+! Identical conditions make the CPU selector a strict superset, zeroing the
+! high-scoring variant. Distinct expressions and declarations must retain the
+! high score even when they fold to the same value.
+! LOCAL-LABEL: func.func @_QMconditionsPlocal_calls(
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPhigh()
+! LOCAL: fir.call @_QMconditionsPhigh()
+! LOCAL: fir.call @_QMconditionsPhigh()
+! LOCAL: fir.call @_QMconditionsPhigh()
+! LOCAL: return
+
+! Conditions referring to USE-associated names are rebuilt in the importing
+! compilation, using the loaded symbols rather than serialized pointers.
+! IMPORTED-LABEL: func.func @_QPimported_conditions(
+! IMPORTED: omp.taskwait
+! IMPORTED: omp.taskwait
+! IMPORTED: omp.taskyield
+! IMPORTED: return
+
+! Same runtime condition: the low-scoring strict superset wins when true.
+! IMPORTED-LABEL: func.func @_QPruntime_same(
+! IMPORTED: fir.if
+! IMPORTED: omp.taskwait
+! IMPORTED: } else {
+! IMPORTED: return
+
+! Different variables must not be merged even when their values might agree.
+! IMPORTED-LABEL: func.func @_QPruntime_distinct(
+! IMPORTED: fir.if
+! IMPORTED: omp.taskyield
+! IMPORTED: } else {
+! IMPORTED: fir.if
+! IMPORTED: omp.taskwait
+! IMPORTED: return
+
+! False conditions retain the same identity under match_none.
+! IMPORTED-LABEL: func.func @_QPfalse_conditions(
+! IMPORTED-NOT: omp.taskyield
+! IMPORTED: omp.taskwait
+! IMPORTED-NOT: omp.taskyield
+! IMPORTED: return
+
+! Both runtime outcomes under match_any retain condition identity, even when
+! the implementation selector makes each candidate applicable independently.
+! IMPORTED-LABEL: func.func @_QPruntime_any(
+! IMPORTED: fir.if
+! IMPORTED: omp.taskwait
+! IMPORTED: } else {
+! IMPORTED: omp.taskwait
+! IMPORTED: return
+
+!--- settings.f90
+module settings
+ logical, parameter :: on = .true., other = .true.
+ logical :: flag, other_flag
+end module
+
+!--- variants.f90
+module conditions
+ use settings, only: on, on2 => on, other
+contains
+ integer function high()
+ high = 100
+ end function
+ integer function low()
+ low = 1
+ end function
+
+ integer function parentheses()
+ !$omp declare variant(high) match(user={condition(.true.)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(((.true.)))}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ parentheses = 0
+ end function
+
+ integer function aliases()
+ !$omp declare variant(high) match(user={condition(on)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(on2)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ aliases = 0
+ end function
+
+ integer function compound_aliases()
+ !$omp declare variant(high) match(user={condition(on .and. .true.)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition((on2.and..true.))}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ compound_aliases = 0
+ end function
+
+ integer function distinct()
+ !$omp declare variant(high) match(user={condition(on)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(other)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ distinct = 0
+ end function
+
+ integer function literal_and_name()
+ !$omp declare variant(high) match(user={condition(.true.)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(on)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ literal_and_name = 0
+ end function
+
+ integer function ordered_expression()
+ !$omp declare variant(high) match(user={condition(on .and. other)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(other .and. on)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ ordered_expression = 0
+ end function
+
+ integer function character_literals()
+ !$omp declare variant(high) match(user={condition('a b' == 'a b')}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition('ab' == 'ab')}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ character_literals = 0
+ end function
+
+ subroutine local_calls(a)
+ integer :: a(7)
+ a(1) = parentheses()
+ a(2) = aliases()
+ a(3) = compound_aliases()
+ a(4) = distinct()
+ a(5) = literal_and_name()
+ a(6) = ordered_expression()
+ a(7) = character_literals()
+ end subroutine
+end module
+
+!--- use.f90
+subroutine imported_conditions()
+ use settings, only: on, alias => on, other
+ !$omp metadirective &
+ !$omp& when(user={condition(on)}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition((alias))}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+ !$omp metadirective &
+ !$omp& when(user={condition(.true.)}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition((.true.))}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+ !$omp metadirective &
+ !$omp& when(user={condition(on)}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition(other)}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+end subroutine
+
+subroutine runtime_same()
+ use settings, only: flag, alias => flag
+ !$omp metadirective &
+ !$omp& when(user={condition(flag)}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition((alias))}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+end subroutine
+
+subroutine runtime_distinct()
+ use settings, only: flag, other_flag
+ !$omp metadirective &
+ !$omp& when(user={condition(flag)}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition(other_flag)}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+end subroutine
+
+subroutine false_conditions()
+ !$omp metadirective &
+ !$omp& when(user={condition(.false.)}, &
+ !$omp& implementation={extension(match_none), vendor(score(100): gnu)}: taskyield) &
+ !$omp& when(user={condition((.false.))}, device={kind(gpu)}, &
+ !$omp& implementation={extension(match_none), vendor(score(1): gnu)}: taskwait)
+end subroutine
+
+subroutine runtime_any()
+ use settings, only: flag, alias => flag
+ !$omp metadirective &
+ !$omp& when(user={condition(flag)}, &
+ !$omp& implementation={extension(match_any), vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition((alias))}, device={kind(cpu)}, &
+ !$omp& implementation={extension(match_any), vendor(score(1): llvm)}: taskwait)
+end subroutine
>From 249d0677013c7049ea5b8dfbb17ee697ebcc1b89 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Sat, 26 Sep 2026 01:53:44 -0500
Subject: [PATCH 08/14] Fix DISTRIBUTE context for SIMD clauses
Enter DISTRIBUTE while evaluating SIMD clauses, then leave it before loop bounds. Cover clause selection, bound exclusion, and restoration.
---
flang/lib/Lower/OpenMP/OpenMP.cpp | 9 ++++-
.../OpenMP/declare-variant-loop-bounds.f90 | 39 +++++++++++++++++++
2 files changed, 46 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index 5e568c1e2965a6..c77f942b4c5e78 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -5532,8 +5532,13 @@ static mlir::omp::DistributeOp genCompositeDistributeSimd(
mlir::omp::SimdOperands simdClauseOps;
llvm::SmallVector<Object> simdReductionObjects;
- genSimdClauses(converter, semaCtx, simdItem->clauses, loc, simdClauseOps,
- simdReductionObjects);
+ {
+ // SIMD clauses see DISTRIBUTE, but loop bounds are evaluated before it.
+ mlir::SaveStateStack<OpenMPContextFrame> distributeContext{
+ converter.getStateStack(), eval, llvm::omp::Directive::OMPD_distribute};
+ genSimdClauses(converter, semaCtx, simdItem->clauses, loc, simdClauseOps,
+ simdReductionObjects);
+ }
DataSharingProcessor distributeItemDSP(
converter, semaCtx, distributeItem->clauses, eval,
diff --git a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90 b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
index 85a301bdfe3a8e..c543c6752a16ef 100644
--- a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
@@ -241,4 +241,43 @@ subroutine nested_simd_clauses(n, a)
end do
end do
end subroutine
+ ! SIMD clauses see TEAMS and DISTRIBUTE (CPU scores 5). The bound and
+ ! post-loop call see only TEAMS (CPU scores 3, losing to vendor score 4).
+ ! CHECK-LABEL: func.func @_QMloop_contextPdistribute_simd_clauses(
+ ! CHECK: omp.teams
+ ! CHECK: fir.call @_QMloop_contextPcpu_count(
+ ! CHECK: fir.call @_QMloop_contextPscored_count(
+ ! CHECK: omp.distribute
+ ! CHECK: omp.simd
+ ! CHECK: fir.call @_QMloop_contextPscored_count(
+ ! CHECK: return
+ subroutine distribute_simd_clauses(n, a)
+ integer :: n, a(n), i
+ !$omp teams
+ !$omp distribute simd if(simd: thread_count(n) > 0)
+ do i = 1, thread_count(n)
+ a(i) = i
+ end do
+ i = thread_count(n)
+ !$omp end teams
+ end subroutine
+
+ ! CHECK-LABEL: func.func @_QMloop_contextPnested_distribute_simd_clauses(
+ ! CHECK: omp.teams
+ ! CHECK: omp.distribute
+ ! CHECK: fir.call @_QMloop_contextPcpu_count(
+ ! CHECK: omp.simd
+ ! CHECK: return
+ subroutine nested_distribute_simd_clauses(n, a)
+ integer :: n, a(n, n), i, j
+ !$omp teams
+ !$omp distribute
+ do i = 1, n
+ !$omp simd if(simd: thread_count(n) > 0)
+ do j = 1, n
+ a(i, j) = i + j
+ end do
+ end do
+ !$omp end teams
+ end subroutine
end module
>From 4328f6d4188999269279a66f9f932258acde9660 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Sat, 26 Sep 2026 01:53:54 -0500
Subject: [PATCH 09/14] Normalize Flang condition expression spelling
Profile original expression structure and declaration identity, normalizing intrinsic operators and literal values and kinds. Preserve signed integer boundaries and cover alternate spellings and nondefault kinds.
---
flang/lib/Semantics/openmp-utils.cpp | 138 ++++++++++-------
.../declare-variant-condition-identity.f90 | 141 +++++++++++++++++-
2 files changed, 220 insertions(+), 59 deletions(-)
diff --git a/flang/lib/Semantics/openmp-utils.cpp b/flang/lib/Semantics/openmp-utils.cpp
index a65b2235445bd5..dac54c1bff2e9c 100644
--- a/flang/lib/Semantics/openmp-utils.cpp
+++ b/flang/lib/Semantics/openmp-utils.cpp
@@ -2485,75 +2485,97 @@ void AppendDirectiveContextTraits(llvm::omp::Directive directive,
}
namespace {
-struct ConditionNames {
- template <typename A> bool Pre(const A &) { return true; }
- template <typename A> void Post(const A &) {}
+// Profile the original parse tree: folding a whole expression would lose
+// declaration identity. Parse-tree operators already unify alternate spellings;
+// analyze only literals to normalize their values and effective kinds.
+struct ConditionIdentity {
+ SemanticsContext &context;
+ llvm::FoldingSetNodeID id;
+
+ template <typename A> void AddNode() {
+ // Type tags and symbol addresses are local to this compilation. Profiles
+ // are rebuilt when importing symbols, never serialized as pointers.
+ static char tag;
+ id.AddPointer(&tag);
+ }
+ template <typename A> bool Pre(const A &x) {
+ AddNode<A>();
+ if constexpr (std::is_enum_v<A> || std::is_integral_v<A>)
+ id.AddInteger(static_cast<uint64_t>(x));
+ else if constexpr (std::is_same_v<A, std::string>)
+ id.AddString(x);
+ return true;
+ }
+ template <typename A> void Post(const A &) { id.AddInteger(0); }
+ // CharBlocks in expression nodes carry source locations, not structure.
+ bool Pre(const parser::CharBlock &) { return false; }
bool Pre(const parser::Name &name) {
+ AddNode<parser::Name>();
if (name.symbol)
- names.push_back(&name);
+ id.AddPointer(&name.symbol->GetUltimate());
+ else
+ id.AddString(name.source.ToString());
+ return false;
+ }
+ bool Pre(const parser::SignedIntLiteralConstant &literal) {
+ AddNode<parser::SignedIntLiteralConstant>();
+ auto value{evaluate::ExpressionAnalyzer{context}.Analyze(literal)};
+ CHECK(value);
+ id.AddString(value->AsFortran());
return false;
}
- llvm::SmallVector<const parser::Name *> names;
+ bool Pre(const parser::SignedRealLiteralConstant &literal) {
+ AddNode<parser::SignedRealLiteralConstant>();
+ auto value{evaluate::ExpressionAnalyzer{context}.Analyze(literal)};
+ CHECK(value);
+ id.AddString(value->AsFortran());
+ return false;
+ }
+ bool Pre(const parser::Expr &expr) {
+ if (const auto *parens{std::get_if<parser::Expr::Parentheses>(&expr.u)}) {
+ parser::Walk(parens->v.value(), *this);
+ return false;
+ }
+ AddNode<parser::Expr>();
+ if (const auto *negate{std::get_if<parser::Expr::Negate>(&expr.u)}) {
+ if (const auto *literal{
+ std::get_if<parser::LiteralConstant>(&negate->v.value().u)}) {
+ if (std::holds_alternative<parser::IntLiteralConstant>(literal->u)) {
+ // The magnitude of the most negative integer is not representable
+ // in its kind. Analyze the signed literal together, as semantics
+ // does, rather than reanalyzing its positive magnitude.
+ auto value{evaluate::ExpressionAnalyzer{context}.Analyze(expr)};
+ CHECK(value);
+ AddNode<parser::Expr::Negate>();
+ id.AddString(value->AsFortran());
+ return false;
+ }
+ }
+ }
+ if (const auto *literal{std::get_if<parser::LiteralConstant>(&expr.u)}) {
+ // Complex literal parts may name declarations: retain their structure.
+ if (!std::holds_alternative<parser::ComplexLiteralConstant>(literal->u)) {
+ if (auto value{evaluate::ExpressionAnalyzer{context}.Analyze(expr)}) {
+ AddNode<parser::LiteralConstant>();
+ id.AddString(value->AsFortran());
+ return false;
+ }
+ }
+ }
+ return true;
+ }
};
-// Profile the original expression, not its folded value. In particular, two
-// named constants with the same value still denote different declarations.
-// Retain the cooked expression spelling, except for outer parentheses, blanks
-// outside character literals, and names resolved to their ultimate symbols.
llvm::StringRef GetConditionIdentity(
const parser::ScalarExpr &condition, SemanticsContext &context) {
const auto *expr{parser::Unwrap<parser::Expr>(condition)};
CHECK(expr);
- while (
- const auto *parentheses{std::get_if<parser::Expr::Parentheses>(&expr->u)})
- expr = &parentheses->v.value();
-
- ConditionNames visitor;
- parser::Walk(*expr, visitor);
- llvm::sort(visitor.names, [](const auto *a, const auto *b) {
- return a->source.begin() < b->source.begin();
- });
-
- llvm::FoldingSetNodeID id;
- auto addText = [&](parser::CharBlock source) {
- std::string text;
- char quote{0};
- for (const char *p{source.begin()}; p != source.end(); ++p) {
- char ch{*p};
- if (quote) {
- text += ch;
- if (ch == quote)
- quote = 0;
- else if (ch == '\\' &&
- context.IsEnabled(common::LanguageFeature::BackslashEscapes) &&
- p + 1 != source.end())
- text += *++p;
- } else if (ch == '\'' || ch == '"') {
- quote = ch;
- text += ch;
- } else if (!llvm::isSpace(ch)) {
- text += ch;
- }
- }
- id.AddString(text);
- };
-
- const char *next{expr->source.begin()};
- for (const parser::Name *name : visitor.names) {
- if (!expr->source.Contains(name->source) || name->source.begin() < next)
- continue;
- addText({next, name->source.begin()});
- id.AddPointer(&name->symbol->GetUltimate());
- next = name->source.end();
- }
- addText({next, expr->source.end()});
-
- // Intern the full profile in context-owned storage. Recompute it after
- // reading a module; symbol pointers are local to this compilation.
- llvm::FoldingSetNodeIDRef profile{id.getRef()};
+ ConditionIdentity profile{context, {}};
+ parser::Walk(*expr, profile);
+ llvm::FoldingSetNodeIDRef data{profile.id.getRef()};
SourceName saved{context.SaveTempName(
- std::string{reinterpret_cast<const char *>(profile.data()),
- profile.size() * sizeof(unsigned)})};
+ std::string{reinterpret_cast<const char *>(data.data()),
+ data.size() * sizeof(unsigned)})};
return {saved.begin(), saved.size()};
}
} // namespace
diff --git a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90 b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
index 93c77029df1338..dd5ac38fe34151 100644
--- a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
@@ -4,6 +4,9 @@
! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir use.f90 -o - | FileCheck %s --check-prefix=IMPORTED
! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -emit-fir use.f90 -o - | FileCheck %s --check-prefix=IMPORTED
+! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -emit-hlfir kinds.f90 -o - | FileCheck %s --check-prefix=DEFAULT4
+! RUN: cd %t && %flang_fc1 -fopenmp -fopenmp-version=52 -fdefault-integer-8 -emit-hlfir kinds.f90 -o - | FileCheck %s --check-prefix=DEFAULT8
+
! Identical conditions make the CPU selector a strict superset, zeroing the
! high-scoring variant. Distinct expressions and declarations must retain the
! high score even when they fold to the same value.
@@ -15,6 +18,16 @@
! LOCAL: fir.call @_QMconditionsPhigh()
! LOCAL: fir.call @_QMconditionsPhigh()
! LOCAL: fir.call @_QMconditionsPhigh()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPhigh()
+! LOCAL: fir.call @_QMconditionsPhigh()
+! LOCAL: fir.call @_QMconditionsPlow()
! LOCAL: return
! Conditions referring to USE-associated names are rebuilt in the importing
@@ -130,8 +143,91 @@ integer function character_literals()
character_literals = 0
end function
+ integer function spelling_eq()
+ !$omp declare variant(high) match(user={condition(1 .eq. 1)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1 == 1)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_eq = 0
+ end function
+
+ integer function spelling_ne()
+ !$omp declare variant(high) match(user={condition(1 .ne. 2)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1 /= 2)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_ne = 0
+ end function
+
+ integer function spelling_lt()
+ !$omp declare variant(high) match(user={condition(1 .lt. 2)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1 < 2)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_lt = 0
+ end function
+
+ integer function spelling_le()
+ !$omp declare variant(high) match(user={condition(1 .le. 1)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1 <= 1)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_le = 0
+ end function
+
+ integer function spelling_gt()
+ !$omp declare variant(high) match(user={condition(2 .gt. 1)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(2 > 1)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_gt = 0
+ end function
+
+ integer function spelling_ge()
+ !$omp declare variant(high) match(user={condition(1 .ge. 1)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1 >= 1)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_ge = 0
+ end function
+
+ integer function spelling_integer_kind()
+ !$omp declare variant(high) match(user={condition(1 == 1)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1_4 == 1_4)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_integer_kind = 0
+ end function
+
+ integer function spelling_distinct_kind()
+ !$omp declare variant(high) match(user={condition(1_4 == 1_4)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1_8 == 1_8)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_distinct_kind = 0
+ end function
+
+ integer function spelling_complex_values()
+ !$omp declare variant(high) match(user={condition((1,2) == (1,2))}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition((1,3) == (1,3))}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_complex_values = 0
+ end function
+
+ ! The positive magnitude is out of range, but the signed literal is valid.
+ integer function spelling_integer_min()
+ !$omp declare variant(high) &
+ !$omp& match(user={condition(-2147483648 == -2147483648)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) &
+ !$omp& match(user={condition(-2147483648_4 == -2147483648_4)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ spelling_integer_min = 0
+ end function
+
subroutine local_calls(a)
- integer :: a(7)
+ integer :: a(17)
a(1) = parentheses()
a(2) = aliases()
a(3) = compound_aliases()
@@ -139,6 +235,16 @@ subroutine local_calls(a)
a(5) = literal_and_name()
a(6) = ordered_expression()
a(7) = character_literals()
+ a(8) = spelling_eq()
+ a(9) = spelling_ne()
+ a(10) = spelling_lt()
+ a(11) = spelling_le()
+ a(12) = spelling_gt()
+ a(13) = spelling_ge()
+ a(14) = spelling_integer_kind()
+ a(15) = spelling_distinct_kind()
+ a(16) = spelling_complex_values()
+ a(17) = spelling_integer_min()
end subroutine
end module
@@ -196,3 +302,36 @@ subroutine runtime_any()
!$omp& when(user={condition((alias))}, device={kind(cpu)}, &
!$omp& implementation={extension(match_any), vendor(score(1): llvm)}: taskwait)
end subroutine
+
+! Imported-symbol runtime expressions also normalize relational operators.
+! IMPORTED-LABEL: func.func @_QPruntime_operators(
+! IMPORTED-NOT: omp.taskyield
+! IMPORTED: fir.if
+! IMPORTED: omp.taskwait
+! IMPORTED-NOT: omp.taskyield
+! IMPORTED: return
+subroutine runtime_operators(n)
+ integer :: n
+ !$omp metadirective &
+ !$omp& when(user={condition(n .eq. 1)}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition(n == 1)}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+end subroutine
+
+!--- kinds.f90
+! DEFAULT4-LABEL: func.func @_QPdefault_kind(
+! DEFAULT4-NOT: omp.taskyield
+! DEFAULT4: omp.taskwait
+! DEFAULT4: return
+! DEFAULT8-LABEL: func.func @_QPdefault_kind(
+! DEFAULT8-NOT: omp.taskwait
+! DEFAULT8: omp.taskyield
+! DEFAULT8: return
+subroutine default_kind()
+ !$omp metadirective &
+ !$omp& when(user={condition(1 == 1)}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition(1_4 == 1_4)}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+end subroutine
>From 4cd0fbd975e5301d00c66d809343a45814d93c82 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Sat, 26 Sep 2026 01:54:07 -0500
Subject: [PATCH 10/14] Preserve metadirective candidate order
Compute subset-adjusted scores once and retain the ordering across runtime guard failures. Share ranks between lowering and reachability so failed supersets cannot restore a subset score.
---
flang/include/flang/Semantics/openmp-utils.h | 12 +--
flang/lib/Lower/OpenMP/OpenMP.cpp | 8 +-
flang/lib/Semantics/openmp-utils.cpp | 78 ++++++++-----------
.../declare-variant-condition-identity.f90 | 8 +-
.../test/Lower/OpenMP/metadirective-user.f90 | 35 +++++++++
.../include/llvm/Frontend/OpenMP/OMPContext.h | 7 ++
llvm/lib/Frontend/OpenMP/OMPContext.cpp | 11 ++-
llvm/unittests/Frontend/OpenMPContextTest.cpp | 19 +++++
8 files changed, 119 insertions(+), 59 deletions(-)
diff --git a/flang/include/flang/Semantics/openmp-utils.h b/flang/include/flang/Semantics/openmp-utils.h
index bdab224907346f..89b0ba55b85b31 100644
--- a/flang/include/flang/Semantics/openmp-utils.h
+++ b/flang/include/flang/Semantics/openmp-utils.h
@@ -265,6 +265,8 @@ struct MetadirectiveCandidate {
bool isExplicit{false};
std::optional<DynamicUserCondition> dynamicCondition;
bool conditionShouldBeTrue{true};
+ /// Position in the original candidate ordering, including subset scores.
+ unsigned rank{0};
};
struct MetadirectiveCandidateSet {
@@ -281,10 +283,10 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet(
const parser::OmpClauseList &clauses, SemanticsContext &context,
const OmpVariantMatchContext &matchContext);
+/// Select the first remaining candidate in the original ranking.
std::optional<unsigned> SelectBestMetadirectiveCandidate(
llvm::ArrayRef<unsigned> candidateIndices,
- llvm::ArrayRef<MetadirectiveCandidate> candidates,
- const OmpVariantMatchContext &matchContext);
+ llvm::ArrayRef<MetadirectiveCandidate> candidates);
/// Return true when repeated evaluation of \p condition cannot call a
/// procedure or observe asynchronously changing state.
@@ -300,13 +302,13 @@ bool AreSameRepeatableMetadirectiveCondition(const parser::ScalarExpr &left,
llvm::SmallVector<unsigned, 4> GetMetadirectiveElsePathCandidates(
unsigned selectedIndex, llvm::ArrayRef<unsigned> candidateIndices,
llvm::ArrayRef<MetadirectiveCandidate> candidates,
- const OmpVariantMatchContext &matchContext, SemanticsContext &context);
+ SemanticsContext &context);
/// Return every replacement that can be selected, retaining lower-ranked
/// candidates after a dynamic condition. Null represents NOTHING.
llvm::SmallVector<const parser::OmpDirectiveSpecification *, 4>
-GetReachableMetadirectiveVariants(const MetadirectiveCandidateSet &candidateSet,
- const OmpVariantMatchContext &matchContext, SemanticsContext &context);
+GetReachableMetadirectiveVariants(
+ const MetadirectiveCandidateSet &candidateSet, SemanticsContext &context);
/// True if a variant guarded by \p selector may be selected in the current
/// compilation context.
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index c77f942b4c5e78..9f334d85faef95 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -7586,7 +7586,7 @@ static void genMetadirective(lower::AbstractConverter &converter,
llvm::SmallVector<const parser::OmpDirectiveSpecification *, 4>
reachableVariantSpecs = semantics::omp::GetReachableMetadirectiveVariants(
- *candidateSet, ompCtx, semaCtx);
+ *candidateSet, semaCtx);
bool hasLoopAssociatedCandidate =
llvm::any_of(reachableVariantSpecs, [](const auto *spec) {
@@ -7823,7 +7823,7 @@ static void genMetadirective(lower::AbstractConverter &converter,
while (!remainingCandidates.empty()) {
std::optional<unsigned> selected =
semantics::omp::SelectBestMetadirectiveCandidate(remainingCandidates,
- candidates, ompCtx);
+ candidates);
if (!selected) {
genVariant(fallback);
return;
@@ -7838,7 +7838,7 @@ static void genMetadirective(lower::AbstractConverter &converter,
llvm::SmallVector<unsigned, 4> elsePathCandidates =
semantics::omp::GetMetadirectiveElsePathCandidates(
- *selected, remainingCandidates, candidates, ompCtx, semaCtx);
+ *selected, remainingCandidates, candidates, semaCtx);
// match_any may create a guarded condition-true candidate and an unguarded
// static candidate for the same directive. If the else path picks the
@@ -7848,7 +7848,7 @@ static void genMetadirective(lower::AbstractConverter &converter,
// else barrier
if (std::optional<unsigned> selectedInElse =
semantics::omp::SelectBestMetadirectiveCandidate(
- elsePathCandidates, candidates, ompCtx)) {
+ elsePathCandidates, candidates)) {
const semantics::omp::MetadirectiveCandidate &candidateInElse =
candidates[*selectedInElse];
if (!candidateInElse.dynamicCondition &&
diff --git a/flang/lib/Semantics/openmp-utils.cpp b/flang/lib/Semantics/openmp-utils.cpp
index dac54c1bff2e9c..d7e48468e1eef0 100644
--- a/flang/lib/Semantics/openmp-utils.cpp
+++ b/flang/lib/Semantics/openmp-utils.cpp
@@ -2809,48 +2809,36 @@ std::optional<MetadirectiveCandidateSet> BuildMetadirectiveCandidateSet(
result.fallback = getFallbackVariant(defaultVariantClause->v.v.value());
}
}
+ // Rank against the complete set once. Removing a failed runtime guard must
+ // not restore the raw score of a selector that was its strict subset.
+ llvm::SmallVector<llvm::omp::VariantMatchInfo, 4> vmis;
+ llvm::SmallVector<unsigned, 4> order;
+ for (const auto &[index, candidate] : llvm::enumerate(result.candidates)) {
+ vmis.push_back(candidate.vmi);
+ order.push_back(index);
+ }
+ auto scores{llvm::omp::getVariantMatchScores(vmis, matchContext)};
+ llvm::stable_sort(order, [&](unsigned a, unsigned b) {
+ CHECK(scores[a] && scores[b]);
+ const auto &left{*scores[a]}, &right{*scores[b]};
+ unsigned width{std::max(left.getBitWidth(), right.getBitWidth())};
+ if (left.zextOrTrunc(width) != right.zextOrTrunc(width))
+ return left.zextOrTrunc(width).ugt(right.zextOrTrunc(width));
+ return result.candidates[a].isExplicit && !result.candidates[b].isExplicit;
+ });
+ for (auto [rank, index] : llvm::enumerate(order))
+ result.candidates[index].rank = rank;
return result;
}
std::optional<unsigned> SelectBestMetadirectiveCandidate(
llvm::ArrayRef<unsigned> candidateIndices,
- llvm::ArrayRef<MetadirectiveCandidate> candidates,
- const OmpVariantMatchContext &matchContext) {
- if (candidateIndices.empty()) {
- return std::nullopt;
- }
- if (candidateIndices.size() == 1) {
- return candidateIndices.front();
- }
-
- // The context scorer preserves input order for ties. Explicit replacements
- // take precedence over an omitted directive's implicit NOTHING.
- llvm::SmallVector<unsigned, 4> candidateOrder;
- candidateOrder.reserve(candidateIndices.size());
- for (unsigned index : candidateIndices) {
- if (candidates[index].isExplicit) {
- candidateOrder.push_back(index);
- }
- }
- for (unsigned index : candidateIndices) {
- if (!candidates[index].isExplicit) {
- candidateOrder.push_back(index);
- }
- }
-
- llvm::SmallVector<llvm::omp::VariantMatchInfo, 4> orderedVMIs;
- orderedVMIs.reserve(candidateOrder.size());
- for (unsigned index : candidateOrder) {
- orderedVMIs.push_back(candidates[index].vmi);
- }
-
- int bestIndex{
- llvm::omp::getBestVariantMatchForContext(orderedVMIs, matchContext)};
- if (bestIndex < 0) {
+ llvm::ArrayRef<MetadirectiveCandidate> candidates) {
+ if (candidateIndices.empty())
return std::nullopt;
- }
- CHECK(static_cast<std::size_t>(bestIndex) < candidateOrder.size());
- return candidateOrder[bestIndex];
+ return *llvm::min_element(candidateIndices, [&](unsigned a, unsigned b) {
+ return candidates[a].rank < candidates[b].rank;
+ });
}
namespace {
@@ -2917,7 +2905,7 @@ bool AreSameRepeatableMetadirectiveCondition(const parser::ScalarExpr &left,
llvm::SmallVector<unsigned, 4> GetMetadirectiveElsePathCandidates(
unsigned selectedIndex, llvm::ArrayRef<unsigned> candidateIndices,
llvm::ArrayRef<MetadirectiveCandidate> candidates,
- const OmpVariantMatchContext &matchContext, SemanticsContext &context) {
+ SemanticsContext &context) {
CHECK(selectedIndex < candidates.size());
const MetadirectiveCandidate &selected{candidates[selectedIndex]};
CHECK(selected.dynamicCondition);
@@ -2934,8 +2922,8 @@ llvm::SmallVector<unsigned, 4> GetMetadirectiveElsePathCandidates(
// condition because it can change state before a lower-ranked occurrence is
// evaluated.
llvm::SmallVector<unsigned, 4> candidatesToInspect{result};
- while (std::optional<unsigned> next{SelectBestMetadirectiveCandidate(
- candidatesToInspect, candidates, matchContext)}) {
+ while (std::optional<unsigned> next{
+ SelectBestMetadirectiveCandidate(candidatesToInspect, candidates)}) {
const MetadirectiveCandidate &candidate{candidates[*next]};
if (!candidate.dynamicCondition ||
!IsRepeatableMetadirectiveCondition(
@@ -2955,8 +2943,8 @@ llvm::SmallVector<unsigned, 4> GetMetadirectiveElsePathCandidates(
}
llvm::SmallVector<const parser::OmpDirectiveSpecification *, 4>
-GetReachableMetadirectiveVariants(const MetadirectiveCandidateSet &candidateSet,
- const OmpVariantMatchContext &matchContext, SemanticsContext &context) {
+GetReachableMetadirectiveVariants(
+ const MetadirectiveCandidateSet &candidateSet, SemanticsContext &context) {
llvm::SmallVector<unsigned, 4> candidates;
candidates.reserve(candidateSet.candidates.size());
for (unsigned index{0}; index < candidateSet.candidates.size(); ++index) {
@@ -2965,8 +2953,8 @@ GetReachableMetadirectiveVariants(const MetadirectiveCandidateSet &candidateSet,
llvm::SmallVector<const parser::OmpDirectiveSpecification *, 4> reachable;
while (true) {
- std::optional<unsigned> selected{SelectBestMetadirectiveCandidate(
- candidates, candidateSet.candidates, matchContext)};
+ std::optional<unsigned> selected{
+ SelectBestMetadirectiveCandidate(candidates, candidateSet.candidates)};
if (!selected) {
reachable.push_back(candidateSet.fallback);
break;
@@ -2981,10 +2969,10 @@ GetReachableMetadirectiveVariants(const MetadirectiveCandidateSet &candidateSet,
}
candidates = GetMetadirectiveElsePathCandidates(
- *selected, candidates, candidateSet.candidates, matchContext, context);
+ *selected, candidates, candidateSet.candidates, context);
if (std::optional<unsigned> selectedInElse{SelectBestMetadirectiveCandidate(
- candidates, candidateSet.candidates, matchContext)}) {
+ candidates, candidateSet.candidates)}) {
const MetadirectiveCandidate &elseCandidate{
candidateSet.candidates[*selectedInElse]};
if (!elseCandidate.dynamicCondition &&
diff --git a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90 b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
index dd5ac38fe34151..145002613f8d35 100644
--- a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
@@ -63,11 +63,13 @@
! Both runtime outcomes under match_any retain condition identity, even when
! the implementation selector makes each candidate applicable independently.
+! Both outcomes select TASKWAIT, so the runtime branch can be eliminated.
! IMPORTED-LABEL: func.func @_QPruntime_any(
-! IMPORTED: fir.if
-! IMPORTED: omp.taskwait
-! IMPORTED: } else {
+! IMPORTED-NOT: fir.if
+! IMPORTED-NOT: omp.taskyield
! IMPORTED: omp.taskwait
+! IMPORTED-NOT: fir.if
+! IMPORTED-NOT: omp.taskyield
! IMPORTED: return
!--- settings.f90
diff --git a/flang/test/Lower/OpenMP/metadirective-user.f90 b/flang/test/Lower/OpenMP/metadirective-user.f90
index 23c92954e7e75d..8c9967a2f2c8d9 100644
--- a/flang/test/Lower/OpenMP/metadirective-user.f90
+++ b/flang/test/Lower/OpenMP/metadirective-user.f90
@@ -678,3 +678,38 @@ subroutine test_multi_dynamic_multi_static(a, b)
!$omp & default(nothing)
#endif
end subroutine
+
+! Rank once: the vendor-only selector scores zero as a subset of the guarded
+! candidate. If that guard fails, its subset must not regain score 101 and
+! overtake the static candidate (score 11). The original order is 22, 11, 0.
+! CHECK-LABEL: func.func @_QPtest_dynamic_subset_order(
+! CHECK: fir.if
+! CHECK-NEXT: omp.barrier
+! CHECK-NEXT: } else {
+! CHECK-NEXT: omp.taskwait
+! CHECK-NOT: omp.taskyield
+! CHECK: return
+subroutine test_dynamic_subset_order(flag)
+ logical :: flag
+ !$omp metadirective &
+ !$omp& when(implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(implementation={vendor(score(1): llvm)}, &
+ !$omp& user={condition(score(20): flag)}: barrier) &
+ !$omp& when(user={condition(score(10): .true.)}: taskwait)
+end subroutine
+
+! The same ordering applies when the guarded replacement is implicit NOTHING.
+! CHECK-LABEL: func.func @_QPtest_dynamic_subset_nothing_order(
+! CHECK: fir.if
+! CHECK-NEXT: } else {
+! CHECK-NEXT: omp.taskwait
+! CHECK-NOT: omp.taskyield
+! CHECK: return
+subroutine test_dynamic_subset_nothing_order(flag)
+ logical :: flag
+ !$omp metadirective &
+ !$omp& when(implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(implementation={vendor(score(1): llvm)}, &
+ !$omp& user={condition(score(20): flag)}:) &
+ !$omp& when(user={condition(score(10): .true.)}: taskwait)
+end subroutine
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h
index f6f3f3a22eeff4..acd0b1bda947d0 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h
@@ -237,6 +237,13 @@ LLVM_ABI bool
isVariantApplicableInContext(const VariantMatchInfo &VMI, const OMPContext &Ctx,
bool DeviceOrImplementationSetOnly = false);
+/// Compute scores for all applicable variants, including strict-subset
+/// adjustments against the complete candidate set. Inapplicable entries have
+/// no score. Keep these scores when testing dynamic candidates in order.
+LLVM_ABI SmallVector<std::optional<APInt>, 4>
+getVariantMatchScores(const SmallVectorImpl<VariantMatchInfo> &VMIs,
+ const OMPContext &Ctx);
+
/// Return the index (into \p VMIs) of the variant with the highest score
/// from the ones applicable in \p Ctx. See llvm::isVariantApplicableInContext.
LLVM_ABI int
diff --git a/llvm/lib/Frontend/OpenMP/OMPContext.cpp b/llvm/lib/Frontend/OpenMP/OMPContext.cpp
index 69d8c94cf98be5..ffa7b7f5e992e4 100644
--- a/llvm/lib/Frontend/OpenMP/OMPContext.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPContext.cpp
@@ -426,8 +426,9 @@ static APInt getVariantMatchScore(const VariantMatchInfo &VMI,
return Score;
}
-int llvm::omp::getBestVariantMatchForContext(
- const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx) {
+SmallVector<std::optional<APInt>, 4>
+llvm::omp::getVariantMatchScores(const SmallVectorImpl<VariantMatchInfo> &VMIs,
+ const OMPContext &Ctx) {
SmallVector<std::optional<APInt>, 4> Scores(VMIs.size());
for (unsigned u = 0, e = VMIs.size(); u < e; ++u) {
const VariantMatchInfo &VMI = VMIs[u];
@@ -455,6 +456,12 @@ int llvm::omp::getBestVariantMatchForContext(
}
}
+ return Scores;
+}
+
+int llvm::omp::getBestVariantMatchForContext(
+ const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx) {
+ auto Scores = getVariantMatchScores(VMIs, Ctx);
APInt BestScore(1, 0);
int BestVMIIdx = -1;
for (unsigned u = 0, e = VMIs.size(); u < e; ++u) {
diff --git a/llvm/unittests/Frontend/OpenMPContextTest.cpp b/llvm/unittests/Frontend/OpenMPContextTest.cpp
index cf8e68f0a33a5b..7a9f95bf6c8872 100644
--- a/llvm/unittests/Frontend/OpenMPContextTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPContextTest.cpp
@@ -574,6 +574,25 @@ TEST_F(OpenMPContextTest, StrictSubsetScoreIsZero) {
EXPECT_EQ(getBestVariantMatchForContext(Candidates, Context), 0);
}
+TEST_F(OpenMPContextTest, PreserveSubsetAdjustedScores) {
+ OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1);
+ APInt High(64, 100), Low(64, 1), Guard(64, 20), Medium(64, 10);
+ VariantMatchInfo Subset, Superset, Other;
+ Subset.addTrait(TraitProperty::implementation_vendor_llvm, "", &High);
+ Superset.addTrait(TraitProperty::implementation_vendor_llvm, "", &Low);
+ Superset.addTrait(TraitProperty::user_condition_true, "flag", &Guard);
+ Other.addTrait(TraitProperty::user_condition_true, "literal", &Medium);
+ SmallVector<VariantMatchInfo, 3> Candidates{Subset, Superset, Other};
+ auto Scores = getVariantMatchScores(Candidates, Context);
+ ASSERT_TRUE(Scores[0] && Scores[1] && Scores[2]);
+ EXPECT_EQ(Scores[0]->getZExtValue(), 0u);
+ EXPECT_EQ(Scores[1]->getZExtValue(), 22u);
+ EXPECT_EQ(Scores[2]->getZExtValue(), 11u);
+ // Consumers can retain these scores after the superset's runtime guard
+ // fails, so Other continues to precede Subset.
+ EXPECT_TRUE(Scores[0]->zext(64).ult(Scores[2]->zext(64)));
+}
+
TEST_F(OpenMPContextTest, IncompatibleSupersetDoesNotZeroScore) {
OMPContext Context(false, Triple("x86_64-unknown-linux"), Triple(), -1);
APInt HighScore(64, 100), LowScore(64, 5);
>From d0f9f012dd2f151af3c7220a639ee6579231b240 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Mon, 28 Sep 2026 11:41:25 -0500
Subject: [PATCH 11/14] Anchor imported condition checks to the selected branch
Require TASKWAIT in the first runtime arm so a nested fallback cannot
satisfy the same-condition expectation.
---
.../test/Lower/OpenMP/declare-variant-condition-identity.f90 | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90 b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
index 145002613f8d35..be0a23d463236f 100644
--- a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
@@ -41,8 +41,8 @@
! Same runtime condition: the low-scoring strict superset wins when true.
! IMPORTED-LABEL: func.func @_QPruntime_same(
! IMPORTED: fir.if
-! IMPORTED: omp.taskwait
-! IMPORTED: } else {
+! IMPORTED-NEXT: omp.taskwait
+! IMPORTED-NEXT: } else {
! IMPORTED: return
! Different variables must not be merged even when their values might agree.
>From 0954321edfbdedb6583510c0f1f5cd5ae80753b6 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Mon, 28 Sep 2026 11:41:26 -0500
Subject: [PATCH 12/14] Use the outer loop context for composite bounds
Resolve original bounds and strides before the outermost loop-associated
constituent. In particular, TEAMS DISTRIBUTE PARALLEL DO bounds see TEAMS,
even when lowering has already created the inner PARALLEL operation.
Scope the context to loop-control evaluation, preserving clause contexts,
genuinely enclosing loops, body contexts, and subsequent expressions.
Cover lower bounds, upper bounds, strides, composite scoring depth, and
static and runtime DO SIMD replacements.
---
flang/lib/Lower/OpenMP/OpenMP.cpp | 4 +-
flang/lib/Lower/OpenMP/Utils.cpp | 30 +++++++
.../OpenMP/declare-variant-loop-bounds.f90 | 79 +++++++++++++++++--
.../declare-variant-metadirective-context.f90 | 46 +++++++++++
4 files changed, 149 insertions(+), 10 deletions(-)
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index 9f334d85faef95..f09b5c8654c3fa 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -7847,8 +7847,8 @@ static void genMetadirective(lower::AbstractConverter &converter,
// if (flag) barrier into just barrier
// else barrier
if (std::optional<unsigned> selectedInElse =
- semantics::omp::SelectBestMetadirectiveCandidate(
- elsePathCandidates, candidates)) {
+ semantics::omp::SelectBestMetadirectiveCandidate(elsePathCandidates,
+ candidates)) {
const semantics::omp::MetadirectiveCandidate &candidateInElse =
candidates[*selectedInElse];
if (!candidateInElse.dynamicCondition &&
diff --git a/flang/lib/Lower/OpenMP/Utils.cpp b/flang/lib/Lower/OpenMP/Utils.cpp
index 8b489b500e8b51..032096c87cae09 100644
--- a/flang/lib/Lower/OpenMP/Utils.cpp
+++ b/flang/lib/Lower/OpenMP/Utils.cpp
@@ -791,6 +791,21 @@ void collectTileSizesFromOpenMPConstruct(
});
}
+namespace {
+// Original loop control belongs outside the outermost loop-associated
+// constituent, including for composites that emit inner regions first.
+class LoopControlContext
+ : public mlir::StateStackFrameBase<LoopControlContext> {
+public:
+ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(LoopControlContext)
+
+ explicit LoopControlContext(const pft::Evaluation &evaluation)
+ : evaluation(evaluation) {}
+
+ const pft::Evaluation &evaluation;
+};
+} // namespace
+
int64_t collectLoopRelatedInfo(
lower::AbstractConverter &converter, mlir::Location currentLocation,
lower::pft::Evaluation &eval, lower::pft::Evaluation *nestedEval,
@@ -823,6 +838,8 @@ void collectLoopRelatedInfo(
int64_t numCollapse, mlir::omp::LoopRelatedClauseOps &result,
llvm::SmallVectorImpl<const semantics::Symbol *> &iv) {
+ mlir::SaveStateStack<LoopControlContext> context{converter.getStateStack(),
+ eval};
fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
// Collect the loops to collapse.
@@ -1521,9 +1538,22 @@ void collectEnclosingConstructTraits(
// PARALLEL when lowering the bounds of PARALLEL DO. A selected directive
// contributes here only through its entered constituents, so its own clause
// expressions have the same context as a directly written directive's.
+ const auto *loopControl =
+ converter.getStateStack().getStackTop<LoopControlContext>();
+ bool insideLoop = false;
for (auto [index, frame] : llvm::enumerate(frames)) {
if (usedFrames[index] || frame->isReplacement)
continue;
+ if (loopControl && &frame->evaluation == &loopControl->evaluation) {
+ // Keep the prefix before the first loop-associated constituent. For
+ // TEAMS DISTRIBUTE PARALLEL DO this is TEAMS, even though the emitted
+ // PARALLEL region already surrounds the bound calculation. Frames for
+ // genuinely enclosing directives are unaffected.
+ insideLoop |= llvm::omp::getDirectiveAssociation(frame->directive) ==
+ llvm::omp::Association::LoopNest;
+ if (insideLoop)
+ continue;
+ }
append(frame->directive);
}
}
diff --git a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90 b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
index c543c6752a16ef..aaa5b4ba5914c5 100644
--- a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
@@ -69,13 +69,14 @@ subroutine explicit_nest(n, a)
!$omp end parallel
end subroutine
- ! Composite lowering creates PARALLEL separately from the loop wrappers.
+ ! Original bounds are evaluated in the context before DISTRIBUTE.
+ ! Emitting the composite PARALLEL region first must not affect selection.
! 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: fir.call @_QMloop_contextPteams_bound(
! CHECK: omp.distribute
! CHECK: omp.wsloop
! CHECK: fir.call @_QMloop_contextPdo_bound(
@@ -95,7 +96,7 @@ subroutine composite(n, a)
! CHECK: omp.parallel
! CHECK-NOT: fir.call @_QMloop_contextPbound(
! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
- ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: fir.call @_QMloop_contextPteams_bound(
! CHECK: omp.distribute
! CHECK: omp.wsloop
! CHECK: omp.simd
@@ -164,13 +165,15 @@ pure integer function thread_count(n)
! DIST_SCHEDULE sees only TEAMS. NUM_THREADS sees TEAMS, DISTRIBUTE, so
! CPU scores 5 and beats the vendor score of 4. A leaked PARALLEL would
- ! make parallel_count a strict superset of cpu_count. The bound sees PARALLEL.
+ ! make parallel_count a strict superset of cpu_count. Bounds see only TEAMS.
! CHECK-LABEL: func.func @_QMloop_contextPcomposite_clauses(
! CHECK: omp.teams
! CHECK: fir.call @_QMloop_contextPteams_bound(
! CHECK: fir.call @_QMloop_contextPcpu_count(
! CHECK: omp.parallel
- ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: fir.call @_QMloop_contextPteams_bound(
+ ! CHECK: fir.call @_QMloop_contextPteams_bound(
+ ! CHECK: fir.call @_QMloop_contextPteams_bound(
! CHECK: omp.distribute
! CHECK: omp.wsloop
! CHECK: fir.call @_QMloop_contextPdo_bound(
@@ -179,7 +182,7 @@ subroutine composite_clauses(n, a)
integer :: n, a(n), i
!$omp teams distribute parallel do dist_schedule(static, bound(n)) &
!$omp& num_threads(thread_count(n))
- do i = 1, bound(n)
+ do i = bound(1), bound(n), bound(1)
a(i) = bound(n)
end do
end subroutine
@@ -191,7 +194,7 @@ subroutine composite_clauses(n, a)
! CHECK: fir.call @_QMloop_contextPcpu_count(
! CHECK: omp.parallel
! CHECK: fir.call @_QMloop_contextPdo_pred(
- ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: fir.call @_QMloop_contextPscored_count(
! CHECK: omp.distribute
! CHECK: omp.wsloop
! CHECK: omp.simd
@@ -201,7 +204,7 @@ subroutine composite_simd_clauses(n, a)
integer :: n, a(n), i
!$omp teams distribute parallel do simd dist_schedule(static, bound(n)) &
!$omp& num_threads(thread_count(n)) if(simd: pred(n))
- do i = 1, bound(n)
+ do i = 1, thread_count(n)
a(i) = bound(n)
end do
end subroutine
@@ -280,4 +283,64 @@ subroutine nested_distribute_simd_clauses(n, a)
end do
!$omp end teams
end subroutine
+
+ ! The same original-bound rule applies without PARALLEL in the composite.
+ ! CHECK-LABEL: func.func @_QMloop_contextPteams_distribute(
+ ! CHECK: omp.teams
+ ! CHECK: fir.call @_QMloop_contextPteams_bound(
+ ! CHECK: fir.call @_QMloop_contextPscored_count(
+ ! CHECK: omp.distribute
+ ! CHECK: fir.call @_QMloop_contextPcpu_count(
+ ! CHECK: fir.call @_QMloop_contextPscored_count(
+ ! CHECK: return
+ subroutine teams_distribute(n, a)
+ integer :: n, a(n), i
+ !$omp teams
+ !$omp distribute
+ do i = bound(1), thread_count(n)
+ a(i) = thread_count(n)
+ end do
+ i = thread_count(n)
+ !$omp end teams
+ end subroutine
+
+ ! Retain PARALLEL preceding DO in the same combined directive.
+ ! CHECK-LABEL: func.func @_QMloop_contextPparallel_do_simd(
+ ! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMloop_contextPdo_pred(
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.simd
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: return
+ subroutine parallel_do_simd(n, a)
+ integer :: n, a(n), i
+ !$omp parallel do simd if(simd: pred(n))
+ do i = 1, bound(n)
+ a(i) = bound(n)
+ end do
+ i = bound(n)
+ end subroutine
+
+ ! An enclosing DO belongs to a different source loop and must remain active
+ ! for the inner SIMD bounds. Only the inner directive's frames are filtered.
+ ! CHECK-LABEL: func.func @_QMloop_contextPenclosing_do(
+ ! CHECK: omp.wsloop
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: omp.simd
+ ! CHECK: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: return
+ subroutine enclosing_do(n, a)
+ integer :: n, a(n,n), i, j
+ !$omp do
+ do i = 1, n
+ !$omp simd
+ do j = 1, bound(n)
+ a(i,j) = bound(n)
+ end do
+ end do
+ i = bound(n)
+ end subroutine
end module
diff --git a/flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90 b/flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90
index 2e5011167c9ee9..b4a30024712967 100644
--- a/flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-metadirective-context.f90
@@ -146,4 +146,50 @@ subroutine delimited_do(n, a)
!$omp end metadirective
a(n) = value(n)
end subroutine
+
+ ! SIMD clauses see DO; original bounds exclude both loop constituents.
+ ! CHECK-LABEL: func.func @_QMselected_contextPstatic_do_simd(
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.simd
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine static_do_simd(n, a)
+ integer :: n, a(n), i
+ !$omp begin metadirective &
+ !$omp& when(implementation={vendor(llvm)}: do simd if(simd: value(n) > 0)) &
+ !$omp& otherwise(nothing)
+ do i = 1, value(n)
+ a(i) = value(n)
+ end do
+ !$omp end metadirective
+ a(n) = value(n)
+ end subroutine
+
+ ! Both runtime arms and the subsequent call restore the outer context.
+ ! CHECK-LABEL: func.func @_QMselected_contextPruntime_do_simd(
+ ! CHECK: fir.if
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: omp.wsloop
+ ! CHECK: omp.simd
+ ! CHECK: fir.call @_QMselected_contextPin_do(
+ ! CHECK: } else {
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: fir.call @_QMselected_contextPvalue(
+ ! CHECK: return
+ subroutine runtime_do_simd(flag, n, a)
+ logical :: flag
+ integer :: n, a(n), i
+ !$omp metadirective &
+ !$omp& when(user={condition(flag)}: do simd if(simd: value(n) > 0)) &
+ !$omp& otherwise(nothing)
+ do i = 1, value(n)
+ a(i) = value(n)
+ end do
+ a(n) = value(n)
+ end subroutine
end module
>From 3080eca811de697fcc3174267ac6a6c2560aa790 Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Tue, 29 Sep 2026 12:09:02 -0500
Subject: [PATCH 13/14] Preserve inner condition parentheses and handle invalid
literals
Strip only outer parentheses from condition identities and avoid assertions
when literal analysis fails. Strengthen runtime branch checks and cover
variant selection during lastprivate bound re-evaluation.
---
flang/lib/Semantics/openmp-utils.cpp | 39 ++++++++--------
.../declare-variant-condition-identity.f90 | 44 ++++++++++++++++---
.../OpenMP/declare-variant-loop-bounds.f90 | 20 +++++++++
.../Semantics/OpenMP/metadirective-user.f90 | 20 +++++++++
4 files changed, 100 insertions(+), 23 deletions(-)
diff --git a/flang/lib/Semantics/openmp-utils.cpp b/flang/lib/Semantics/openmp-utils.cpp
index d7e48468e1eef0..a773ec3b7df936 100644
--- a/flang/lib/Semantics/openmp-utils.cpp
+++ b/flang/lib/Semantics/openmp-utils.cpp
@@ -2487,7 +2487,8 @@ void AppendDirectiveContextTraits(llvm::omp::Directive directive,
namespace {
// Profile the original parse tree: folding a whole expression would lose
// declaration identity. Parse-tree operators already unify alternate spellings;
-// analyze only literals to normalize their values and effective kinds.
+// analyze only literals to normalize their values and effective kinds. Failed
+// analysis during semantic error recovery falls back to the structural walk.
struct ConditionIdentity {
SemanticsContext &context;
llvm::FoldingSetNodeID id;
@@ -2519,23 +2520,21 @@ struct ConditionIdentity {
}
bool Pre(const parser::SignedIntLiteralConstant &literal) {
AddNode<parser::SignedIntLiteralConstant>();
- auto value{evaluate::ExpressionAnalyzer{context}.Analyze(literal)};
- CHECK(value);
- id.AddString(value->AsFortran());
- return false;
+ if (auto value{evaluate::ExpressionAnalyzer{context}.Analyze(literal)}) {
+ id.AddString(value->AsFortran());
+ return false;
+ }
+ return true;
}
bool Pre(const parser::SignedRealLiteralConstant &literal) {
AddNode<parser::SignedRealLiteralConstant>();
- auto value{evaluate::ExpressionAnalyzer{context}.Analyze(literal)};
- CHECK(value);
- id.AddString(value->AsFortran());
- return false;
- }
- bool Pre(const parser::Expr &expr) {
- if (const auto *parens{std::get_if<parser::Expr::Parentheses>(&expr.u)}) {
- parser::Walk(parens->v.value(), *this);
+ if (auto value{evaluate::ExpressionAnalyzer{context}.Analyze(literal)}) {
+ id.AddString(value->AsFortran());
return false;
}
+ return true;
+ }
+ bool Pre(const parser::Expr &expr) {
AddNode<parser::Expr>();
if (const auto *negate{std::get_if<parser::Expr::Negate>(&expr.u)}) {
if (const auto *literal{
@@ -2544,11 +2543,11 @@ struct ConditionIdentity {
// The magnitude of the most negative integer is not representable
// in its kind. Analyze the signed literal together, as semantics
// does, rather than reanalyzing its positive magnitude.
- auto value{evaluate::ExpressionAnalyzer{context}.Analyze(expr)};
- CHECK(value);
- AddNode<parser::Expr::Negate>();
- id.AddString(value->AsFortran());
- return false;
+ if (auto value{evaluate::ExpressionAnalyzer{context}.Analyze(expr)}) {
+ AddNode<parser::Expr::Negate>();
+ id.AddString(value->AsFortran());
+ return false;
+ }
}
}
}
@@ -2570,6 +2569,10 @@ llvm::StringRef GetConditionIdentity(
const parser::ScalarExpr &condition, SemanticsContext &context) {
const auto *expr{parser::Unwrap<parser::Expr>(condition)};
CHECK(expr);
+ // Only outer parentheses are insignificant. Inner parentheses can constrain
+ // reassociation and must remain part of the original expression's identity.
+ while (const auto *parens{std::get_if<parser::Expr::Parentheses>(&expr->u)})
+ expr = &parens->v.value();
ConditionIdentity profile{context, {}};
parser::Walk(*expr, profile);
llvm::FoldingSetNodeIDRef data{profile.id.getRef()};
diff --git a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90 b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
index be0a23d463236f..d890786755f51b 100644
--- a/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-condition-identity.f90
@@ -28,6 +28,7 @@
! LOCAL: fir.call @_QMconditionsPhigh()
! LOCAL: fir.call @_QMconditionsPhigh()
! LOCAL: fir.call @_QMconditionsPlow()
+! LOCAL: fir.call @_QMconditionsPhigh()
! LOCAL: return
! Conditions referring to USE-associated names are rebuilt in the importing
@@ -48,10 +49,10 @@
! Different variables must not be merged even when their values might agree.
! IMPORTED-LABEL: func.func @_QPruntime_distinct(
! IMPORTED: fir.if
-! IMPORTED: omp.taskyield
-! IMPORTED: } else {
+! IMPORTED-NEXT: omp.taskyield
+! IMPORTED-NEXT: } else {
! IMPORTED: fir.if
-! IMPORTED: omp.taskwait
+! IMPORTED-NEXT: omp.taskwait
! IMPORTED: return
! False conditions retain the same identity under match_none.
@@ -228,8 +229,17 @@ integer function spelling_integer_min()
spelling_integer_min = 0
end function
+ ! Inner parentheses remain significant even if both expressions fold alike.
+ integer function inner_parentheses()
+ !$omp declare variant(high) match(user={condition((1)+1 == 2)}, &
+ !$omp& implementation={vendor(score(100): llvm)})
+ !$omp declare variant(low) match(user={condition(1+1 == 2)}, &
+ !$omp& implementation={vendor(score(1): llvm)}, device={kind(cpu)})
+ inner_parentheses = 0
+ end function
+
subroutine local_calls(a)
- integer :: a(17)
+ integer :: a(18)
a(1) = parentheses()
a(2) = aliases()
a(3) = compound_aliases()
@@ -247,6 +257,7 @@ subroutine local_calls(a)
a(15) = spelling_distinct_kind()
a(16) = spelling_complex_values()
a(17) = spelling_integer_min()
+ a(18) = inner_parentheses()
end subroutine
end module
@@ -309,7 +320,8 @@ subroutine runtime_any()
! IMPORTED-LABEL: func.func @_QPruntime_operators(
! IMPORTED-NOT: omp.taskyield
! IMPORTED: fir.if
-! IMPORTED: omp.taskwait
+! IMPORTED-NEXT: omp.taskwait
+! IMPORTED-NEXT: } else {
! IMPORTED-NOT: omp.taskyield
! IMPORTED: return
subroutine runtime_operators(n)
@@ -321,6 +333,28 @@ subroutine runtime_operators(n)
!$omp& implementation={vendor(score(1): llvm)}: taskwait)
end subroutine
+! Inner parentheses constrain reassociation. The higher-scoring condition
+! must be tested first, while an outer pair around it remains insignificant.
+! IMPORTED-LABEL: func.func @_QPruntime_inner_parentheses(
+! IMPORTED: %[[SUM:.*]] = arith.addf
+! IMPORTED: %[[GROUP:.*]] = {{(hlfir|fir)}}.no_reassoc %[[SUM]]
+! IMPORTED: arith.addf %[[GROUP]],
+! IMPORTED: fir.if
+! IMPORTED-NEXT: omp.taskyield
+! IMPORTED-NEXT: } else {
+! IMPORTED: fir.if
+! IMPORTED-NEXT: omp.taskwait
+! IMPORTED-NEXT: } else {
+! IMPORTED: return
+subroutine runtime_inner_parentheses(a, b, c)
+ real :: a, b, c
+ !$omp metadirective &
+ !$omp& when(user={condition(((a+b)+c > 0.0))}, &
+ !$omp& implementation={vendor(score(100): llvm)}: taskyield) &
+ !$omp& when(user={condition(a+b+c > 0.0)}, device={kind(cpu)}, &
+ !$omp& implementation={vendor(score(1): llvm)}: taskwait)
+end subroutine
+
!--- kinds.f90
! DEFAULT4-LABEL: func.func @_QPdefault_kind(
! DEFAULT4-NOT: omp.taskyield
diff --git a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90 b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
index aaa5b4ba5914c5..bff53d8ade266d 100644
--- a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
@@ -343,4 +343,24 @@ subroutine enclosing_do(n, a)
end do
i = bound(n)
end subroutine
+ ! Lastprivate re-evaluates the bound inside WSLOOP for the copy-back test.
+ ! Both evaluations belong to the original loop control, outside DO.
+ ! CHECK-LABEL: func.func @_QMloop_contextPlastprivate_bound(
+ ! CHECK: fir.call @_QMloop_contextPbound(
+ ! CHECK: omp.wsloop
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: %[[COPY_UB:.*]] = fir.call @_QMloop_contextPbound(
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: arith.cmpi sgt, %{{[^,]+}}, %[[COPY_UB]]
+ ! CHECK: fir.if
+ ! CHECK: {{hlfir.assign|fir.store}}
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine lastprivate_bound(n, x)
+ integer :: n, x, i
+ !$omp do lastprivate(x)
+ do i = 1, bound(n)
+ x = i
+ end do
+ end subroutine
end module
diff --git a/flang/test/Semantics/OpenMP/metadirective-user.f90 b/flang/test/Semantics/OpenMP/metadirective-user.f90
index 202d7781cd7718..12e2dd0f028cf4 100644
--- a/flang/test/Semantics/OpenMP/metadirective-user.f90
+++ b/flang/test/Semantics/OpenMP/metadirective-user.f90
@@ -34,3 +34,23 @@ subroutine f04(a, b)
!ERROR: Repeated trait name CONDITION in a trait set
!$omp & when(user={condition(a), condition(b)}: nothing)
end
+
+! Profiling still runs after expression diagnostics have been recorded.
+subroutine invalid_negative_kind(n)
+ integer :: n
+!ERROR: CONDITION trait requires a single LOGICAL expression
+!ERROR: INTEGER(KIND=3) is not a supported type
+ !$omp metadirective when(user={condition(n == -1_3)}: taskyield) otherwise(taskwait)
+end
+
+subroutine invalid_complex_integer_kind()
+!ERROR: CONDITION trait requires a single LOGICAL expression
+!ERROR: INTEGER(KIND=3) is not a supported type
+ !$omp metadirective when(user={condition((1_3, 2) == (1, 2))}: taskyield) otherwise(taskwait)
+end
+
+subroutine invalid_complex_real_kind()
+!ERROR: CONDITION trait requires a single LOGICAL expression
+!ERROR: Unsupported REAL(KIND=7)
+ !$omp metadirective when(user={condition((1.0_7, 0.0) == (1.0, 0.0))}: taskyield) otherwise(taskwait)
+end
>From 89a879d2231f0ef2983318759a6354ec279df29d Mon Sep 17 00:00:00 2001
From: "Chi-Chun, Chen" <chichun.chen at hpe.com>
Date: Wed, 30 Sep 2026 16:55:57 -0500
Subject: [PATCH 14/14] Fix variant context for lastprivate bounds
Use the owning directive's evaluation when collecting construct ancestors
for loop-control expressions. Lastprivate can re-evaluate bounds after
loop-body lowering leaves a body evaluation current, which otherwise adds
the loop construct before the loop-control context filter runs.
---
flang/lib/Lower/OpenMP/Utils.cpp | 10 ++++-
.../OpenMP/declare-variant-loop-bounds.f90 | 43 +++++++++++++++++++
2 files changed, 51 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Lower/OpenMP/Utils.cpp b/flang/lib/Lower/OpenMP/Utils.cpp
index 032096c87cae09..346635375ab78f 100644
--- a/flang/lib/Lower/OpenMP/Utils.cpp
+++ b/flang/lib/Lower/OpenMP/Utils.cpp
@@ -1492,6 +1492,14 @@ semantics::omp::OmpVariantMatchContext makeVariantMatchContext(
void collectEnclosingConstructTraits(
AbstractConverter &converter, const pft::Evaluation *evaluation,
llvm::SmallVectorImpl<llvm::omp::TraitProperty> &constructTraits) {
+ const auto *loopControl =
+ converter.getStateStack().getStackTop<LoopControlContext>();
+ // Lastprivate can re-evaluate bounds after lowering the loop body, leaving
+ // a body evaluation current. Use the owning directive's ancestors so the
+ // loop itself is not added before filtering its entered constituents below.
+ if (loopControl)
+ evaluation = &loopControl->evaluation;
+
llvm::SmallVector<const OpenMPContextFrame *, 4> frames;
converter.getStateStack().stackWalk<OpenMPContextFrame>(
[&](OpenMPContextFrame &frame) {
@@ -1538,8 +1546,6 @@ void collectEnclosingConstructTraits(
// PARALLEL when lowering the bounds of PARALLEL DO. A selected directive
// contributes here only through its entered constituents, so its own clause
// expressions have the same context as a directly written directive's.
- const auto *loopControl =
- converter.getStateStack().getStackTop<LoopControlContext>();
bool insideLoop = false;
for (auto [index, frame] : llvm::enumerate(frames)) {
if (usedFrames[index] || frame->isReplacement)
diff --git a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90 b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
index bff53d8ade266d..de2e13d58973e5 100644
--- a/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
+++ b/flang/test/Lower/OpenMP/declare-variant-loop-bounds.f90
@@ -363,4 +363,47 @@ subroutine lastprivate_bound(n, x)
x = i
end do
end subroutine
+
+ ! Keep the entered PARALLEL constituent when re-evaluating the bound.
+ ! CHECK-LABEL: func.func @_QMloop_contextPlastprivate_combined(
+ ! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.wsloop
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: %[[PAR_UB:.*]] = fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: arith.cmpi sgt, %{{[^,]+}}, %[[PAR_UB]]
+ ! CHECK: fir.if
+ ! CHECK: {{hlfir.assign|fir.store}}
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine lastprivate_combined(n, x)
+ integer :: n, x, i
+ !$omp parallel do lastprivate(x)
+ do i = 1, bound(n)
+ x = i
+ end do
+ end subroutine
+
+ ! A separately enclosing PARALLEL must also remain in the bound context.
+ ! CHECK-LABEL: func.func @_QMloop_contextPlastprivate_nested(
+ ! CHECK: omp.parallel
+ ! CHECK: fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: omp.wsloop
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: %[[NESTED_UB:.*]] = fir.call @_QMloop_contextPparallel_bound(
+ ! CHECK: arith.cmpi sgt, %{{[^,]+}}, %[[NESTED_UB]]
+ ! CHECK: fir.if
+ ! CHECK: {{hlfir.assign|fir.store}}
+ ! CHECK-NOT: fir.call @_QMloop_contextPdo_bound(
+ ! CHECK: return
+ subroutine lastprivate_nested(n, x)
+ integer :: n, x, i
+ !$omp parallel
+ !$omp do lastprivate(x)
+ do i = 1, bound(n)
+ x = i
+ end do
+ !$omp end parallel
+ end subroutine
+
end module
More information about the llvm-commits
mailing list