[clang] [OpenMP] Map const-qualified target map variables as 'to'. (PR #185918)
Zahira Ammarguellat via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 12 07:21:29 PDT 2026
https://github.com/zahiraam updated https://github.com/llvm/llvm-project/pull/185918
>From 310f7226b1dbc101823253d99103b4a300f73635 Mon Sep 17 00:00:00 2001
From: Ammarguellat <zahira.ammarguellat at intel.com>
Date: Wed, 11 Mar 2026 09:38:48 -0700
Subject: [PATCH 1/3] [OpenMP] Map const-qualified target map variables as
'to'.
---
clang/lib/Sema/SemaOpenMP.cpp | 38 +++++-
clang/test/OpenMP/map_const_aggregate.cpp | 150 ++++++++++++++++++++++
2 files changed, 184 insertions(+), 4 deletions(-)
create mode 100644 clang/test/OpenMP/map_const_aggregate.cpp
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index e8a7c1d5e9288..5ff6b9978cdd8 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -3862,7 +3862,8 @@ static void reportOriginalDsa(Sema &SemaRef, const DSAStackTy *Stack,
static OpenMPMapClauseKind
getMapClauseKindFromModifier(OpenMPDefaultmapClauseModifier M,
- bool IsAggregateOrDeclareTarget) {
+ bool IsAggregateOrDeclareTarget,
+ bool HasConstQualifier) {
OpenMPMapClauseKind Kind = OMPC_MAP_unknown;
switch (M) {
case OMPC_DEFAULTMAP_MODIFIER_alloc:
@@ -3897,7 +3898,10 @@ getMapClauseKindFromModifier(OpenMPDefaultmapClauseModifier M,
// 1. the implicit behavior for aggregate is tofrom
// 2. it's a declare target link
if (IsAggregateOrDeclareTarget) {
- Kind = OMPC_MAP_tofrom;
+ if (HasConstQualifier)
+ Kind = OMPC_MAP_to;
+ else
+ Kind = OMPC_MAP_tofrom;
break;
}
llvm_unreachable("Unexpected defaultmap implicit behavior");
@@ -3906,6 +3910,30 @@ getMapClauseKindFromModifier(OpenMPDefaultmapClauseModifier M,
return Kind;
}
+static bool hasNoMutableFields(const CXXRecordDecl *RD) {
+ for (const auto *FD : RD->fields()) {
+ if (FD->isMutable())
+ return false;
+ QualType FT = FD->getType();
+ while (FT->isArrayType())
+ FT = FT->getAsArrayTypeUnsafe()->getElementType();
+ if (const auto *NestedRD = FT->getAsCXXRecordDecl())
+ if (!hasNoMutableFields(NestedRD))
+ return false;
+ }
+ return true;
+}
+
+static bool hasConstQualifiedMappingType(QualType T) {
+ while (T->isArrayType())
+ T = T->getAsArrayTypeUnsafe()->getElementType();
+ if (!T.isConstQualified())
+ return false;
+ if (const auto *RD = T->getAsCXXRecordDecl())
+ return hasNoMutableFields(RD);
+ return true;
+}
+
namespace {
struct VariableImplicitInfo {
static const unsigned MapKindNum = OMPC_MAP_unknown;
@@ -4128,7 +4156,8 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
ImpInfo.Privates.insert(E);
} else {
OpenMPMapClauseKind Kind = getMapClauseKindFromModifier(
- M, ClauseKind == OMPC_DEFAULTMAP_aggregate || Res);
+ M, ClauseKind == OMPC_DEFAULTMAP_aggregate || Res,
+ hasConstQualifiedMappingType(E->getType()));
ImpInfo.Mappings[ClauseKind][Kind].insert(E);
}
}
@@ -4225,7 +4254,8 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
OpenMPDefaultmapClauseKind ClauseKind =
getVariableCategoryFromDecl(SemaRef.getLangOpts(), FD);
OpenMPMapClauseKind Kind = getMapClauseKindFromModifier(
- Modifier, /*IsAggregateOrDeclareTarget=*/true);
+ Modifier, /*IsAggregateOrDeclareTarget=*/true,
+ /*HasConstQualifier=*/false);
ImpInfo.Mappings[ClauseKind][Kind].insert(E);
return;
}
diff --git a/clang/test/OpenMP/map_const_aggregate.cpp b/clang/test/OpenMP/map_const_aggregate.cpp
new file mode 100644
index 0000000000000..857a525533ca2
--- /dev/null
+++ b/clang/test/OpenMP/map_const_aggregate.cpp
@@ -0,0 +1,150 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu \
+// RUN: -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+
+// RUN %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu \
+// RUN -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+
+// expected-no-diagnostics
+
+struct foo {
+ foo(int j) : i(j) {};
+ int i;
+};
+
+// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 545]
+// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 547]
+// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 547]
+// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [1 x i64] [i64 545]
+// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [1 x i64] [i64 547]
+// CHECK: @.offload_maptypes.10 = private unnamed_addr constant [1 x i64] [i64 545]
+// CHECK: @.offload_maptypes.12 = private unnamed_addr constant [1 x i64] [i64 35]
+// CHECK: @.offload_maptypes.14 = private unnamed_addr constant [2 x i64] [i64 545, i64 547]
+
+// Const struct, no mutable members -> mapped 'to'
+
+// LABEL: test_const_no_mutable
+// CHECK: store ptr @.offload_maptypes, ptr {{.*}}, align 8
+void test_const_no_mutable() {
+ const foo a(2);
+#pragma omp target
+ {
+ int x = a.i;
+ }
+}
+
+// Non-const -> mapped 'tofrom'
+
+// LABEL: define dso_local void @_Z13test_nonconstv
+// CHECK: store ptr @.offload_maptypes.2, ptr {{.*}}, align 8
+void test_nonconst() {
+ foo a(2);
+#pragma omp target
+ {
+ int x = a.i;
+ }
+}
+
+struct foo_mutable {
+ foo_mutable(int j) : i(j), m(0) {};
+ int i;
+ mutable int m;
+};
+
+// Const struct with a mutable member -> mapped 'tofrom'
+
+// LABEL: define dso_local void @_Z23test_const_with_mutablev
+// CHECK: store ptr @.offload_maptypes.4, ptr {{.*}}, align 8
+void test_const_with_mutable() {
+ const foo_mutable a(2);
+#pragma omp target
+ {
+ a.m = 1;
+ }
+}
+
+struct foo_nested {
+ foo_nested(int j) : inner(j), z(j) {};
+ foo inner;
+ const int z;
+};
+
+// Const struct nested inside another const struct -> mapped 'to'
+
+// LABEL: define dso_local void @_Z17test_const_nestedv() #0 {
+// CHECK: store ptr @.offload_maptypes.6, ptr {{.*}}, align 8
+void test_const_nested() {
+ const foo_nested a(2);
+#pragma omp target
+ {
+ int x = a.inner.i;
+ }
+}
+
+struct foo_nested_mutable {
+ foo_nested_mutable(int j) : inner(j), z(j) {};
+ foo_mutable inner; // has mutable member buried inside
+ const int z;
+};
+
+// Const struct nested inside another const struct, where the nested
+// struct has a mutable member -> mapped 'tofrom'
+
+// LABEL: define dso_local void @_Z30test_const_nested_with_mutablev
+// CHECK: store ptr @.offload_maptypes.8, ptr {{.*}}, align 8
+void test_const_nested_with_mutable() {
+ const foo_nested_mutable a(2);
+#pragma omp target
+ {
+ a.inner.m = 1;
+ }
+}
+
+// Const array of foo -> mapped 'to'
+
+// LABEL: define dso_local void @_Z16test_const_arrayv
+// CHECK: store ptr @.offload_maptypes.10, ptr {{.*}}, align 8
+void test_const_array() {
+ const foo arr[4] = {1, 2, 3, 4};
+#pragma omp target
+ {
+ int x = arr[0].i;
+ }
+}
+
+// Explicit map(tofrom:) on a const struct -> mapped 'tofrom'
+
+// LABEL: define dso_local void @_Z27test_explicit_map_overridesv
+// CHECK: store ptr @.offload_maptypes.12, ptr {{.*}}, align 8
+void test_explicit_map_overrides() {
+ const foo a(2);
+#pragma omp target map(tofrom:a)
+ {
+ int x = a.i;
+ }
+}
+
+// Mixed: const foo (to) and non-const foo (tofrom) in the same region.
+
+// LABEL: define dso_local void @_Z10test_mixedv
+// CHECK: store ptr @.offload_maptypes.14, ptr {{.*}}, align 8
+void test_mixed() {
+ const foo ca(2);
+ foo ma(3);
+#pragma omp target
+ {
+ int x = ca.i;
+ ma.i = 99;
+ }
+}
+
+// Defaultmap(tofrom:aggregate) explicit -> mapped 'to'.
+
+// LABEL: define dso_local void @_Z31test_defaultmap_tofrom_explicitv
+// CHECK: store ptr @.offload_maptypes.16, ptr {{.*}}, align 8
+void test_defaultmap_tofrom_explicit() {
+ const foo a(2);
+#pragma omp target defaultmap(tofrom:aggregate)
+ {
+ int x = a.i;
+ }
+}
>From 9ea784f31952828283c4aadc9d819bcaa2ba3ab8 Mon Sep 17 00:00:00 2001
From: Ammarguellat <zahira.ammarguellat at intel.com>
Date: Wed, 11 Mar 2026 09:47:01 -0700
Subject: [PATCH 2/3] Fix format
---
clang/lib/Sema/SemaOpenMP.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 5ff6b9978cdd8..4a1003a406b35 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -3899,7 +3899,7 @@ getMapClauseKindFromModifier(OpenMPDefaultmapClauseModifier M,
// 2. it's a declare target link
if (IsAggregateOrDeclareTarget) {
if (HasConstQualifier)
- Kind = OMPC_MAP_to;
+ Kind = OMPC_MAP_to;
else
Kind = OMPC_MAP_tofrom;
break;
>From 89eb6a0d97e4a8643e596c919bf18bb0ca16fe14 Mon Sep 17 00:00:00 2001
From: Ammarguellat <zahira.ammarguellat at intel.com>
Date: Thu, 12 Mar 2026 07:21:14 -0700
Subject: [PATCH 3/3] Addressed review comments
---
clang/lib/Sema/SemaOpenMP.cpp | 10 ++
clang/test/OpenMP/map_const_aggregate.cpp | 157 ++++++++++++++++------
2 files changed, 125 insertions(+), 42 deletions(-)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 4a1003a406b35..05b6e31751a96 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -23510,6 +23510,16 @@ OMPClause *SemaOpenMP::ActOnOpenMPMapClause(
}
MappableVarListInfo MVLI(VarList);
+ // Per OpenMP 6.0 p299 lines 3-4, a list item with the const specifier and
+ // no mutable members is ignored for 'from' clauses. A const-qualified
+ // variable cannot be modified on the device, so copying back to the host
+ // is unnecessary and potentially unsafe. Strip the FROM component:
+ // map(tofrom:) -> map(to:), map(from:) -> map(alloc:).
+ for (auto *E : VarList) {
+ if ((MapType == OMPC_MAP_from || MapType == OMPC_MAP_tofrom) &&
+ hasConstQualifiedMappingType(E->getType()))
+ MapType = (MapType == OMPC_MAP_tofrom) ? OMPC_MAP_to : OMPC_MAP_alloc;
+ }
checkMappableExpressionList(SemaRef, DSAStack, OMPC_map, MVLI, Locs.StartLoc,
MapperIdScopeSpec, MapperId, UnresolvedMappers,
MapType, Modifiers, IsMapTypeImplicit,
diff --git a/clang/test/OpenMP/map_const_aggregate.cpp b/clang/test/OpenMP/map_const_aggregate.cpp
index 857a525533ca2..ed9d67bc4c064 100644
--- a/clang/test/OpenMP/map_const_aggregate.cpp
+++ b/clang/test/OpenMP/map_const_aggregate.cpp
@@ -6,22 +6,54 @@
// expected-no-diagnostics
+// Tests that const-qualified aggregates without mutable members are implicitly
+// mapped as 'to' instead of 'tofrom' under defaultmap(tofrom:aggregate) and
+// explicit map clauses. Structs that have mutable members, or that are
+// non-const, must continue to be mapped 'tofrom'.
+
struct foo {
foo(int j) : i(j) {};
int i;
};
+struct foo_mutable {
+ foo_mutable(int j) : i(j), m(0) {};
+ int i;
+ mutable int m;
+};
+
+struct foo_nested {
+ foo_nested(int j) : inner(j), z(j) {};
+ foo inner;
+ const int z;
+};
+
+struct foo_nested_mutable {
+ foo_nested_mutable(int j) : inner(j), z(j) {};
+ foo_mutable inner; // has mutable member buried inside
+ const int z;
+};
+
// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 545]
// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 547]
// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 547]
// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [1 x i64] [i64 545]
// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [1 x i64] [i64 547]
// CHECK: @.offload_maptypes.10 = private unnamed_addr constant [1 x i64] [i64 545]
-// CHECK: @.offload_maptypes.12 = private unnamed_addr constant [1 x i64] [i64 35]
-// CHECK: @.offload_maptypes.14 = private unnamed_addr constant [2 x i64] [i64 545, i64 547]
-
-// Const struct, no mutable members -> mapped 'to'
-
+// CHECK: @.offload_maptypes.12 = private unnamed_addr constant [1 x i64] [i64 33]
+// CHECK: @.offload_maptypes.14 = private unnamed_addr constant [1 x i64] [i64 32]
+// CHECK: @.offload_maptypes.16 = private unnamed_addr constant [1 x i64] [i64 33]
+// CHECK: @.offload_maptypes.18 = private unnamed_addr constant [1 x i64] [i64 2]
+// CHECK: @.offload_maptypes.20 = private unnamed_addr constant [1 x i64] [i64 2]
+// CHECK: @.offload_maptypes.22 = private unnamed_addr constant [1 x i64] [i64 2]
+// CHECK: @.offload_maptypes.24 = private unnamed_addr constant [2 x i64] [i64 545, i64 547]
+// CHECK: @.offload_maptypes.26 = private unnamed_addr constant [1 x i64] [i64 545]
+
+// ---------------------------------------------------------------------------
+// Implicit mapping tests (no explicit map clause, defaultmap governs)
+// ---------------------------------------------------------------------------
+
+// Const struct with no mutable members. Mapped as TO|TARGET_PARAM|IMPLICIT = 545.
// LABEL: test_const_no_mutable
// CHECK: store ptr @.offload_maptypes, ptr {{.*}}, align 8
void test_const_no_mutable() {
@@ -32,8 +64,7 @@ void test_const_no_mutable() {
}
}
-// Non-const -> mapped 'tofrom'
-
+// Non-const struct. Mapped as TO|FROM|TARGET_PARAM|IMPLICIT = 547.
// LABEL: define dso_local void @_Z13test_nonconstv
// CHECK: store ptr @.offload_maptypes.2, ptr {{.*}}, align 8
void test_nonconst() {
@@ -44,14 +75,7 @@ void test_nonconst() {
}
}
-struct foo_mutable {
- foo_mutable(int j) : i(j), m(0) {};
- int i;
- mutable int m;
-};
-
-// Const struct with a mutable member -> mapped 'tofrom'
-
+// Const struct with a mutable member. Mapped as TO|FROM|TARGET_PARAM|IMPLICIT = 547.
// LABEL: define dso_local void @_Z23test_const_with_mutablev
// CHECK: store ptr @.offload_maptypes.4, ptr {{.*}}, align 8
void test_const_with_mutable() {
@@ -62,14 +86,8 @@ void test_const_with_mutable() {
}
}
-struct foo_nested {
- foo_nested(int j) : inner(j), z(j) {};
- foo inner;
- const int z;
-};
-
-// Const struct nested inside another const struct -> mapped 'to'
-
+// Const struct whose members are themselves all const and free of mutable
+// fields. Mapped as TO|TARGET_PARAM|IMPLICIT = 545.
// LABEL: define dso_local void @_Z17test_const_nestedv() #0 {
// CHECK: store ptr @.offload_maptypes.6, ptr {{.*}}, align 8
void test_const_nested() {
@@ -80,15 +98,8 @@ void test_const_nested() {
}
}
-struct foo_nested_mutable {
- foo_nested_mutable(int j) : inner(j), z(j) {};
- foo_mutable inner; // has mutable member buried inside
- const int z;
-};
-
-// Const struct nested inside another const struct, where the nested
-// struct has a mutable member -> mapped 'tofrom'
-
+// Const array of a const-qualified struct type.
+// Mapped as TO|FROM|TARGET_PARAM|IMPLICIT = 547.
// LABEL: define dso_local void @_Z30test_const_nested_with_mutablev
// CHECK: store ptr @.offload_maptypes.8, ptr {{.*}}, align 8
void test_const_nested_with_mutable() {
@@ -99,8 +110,8 @@ void test_const_nested_with_mutable() {
}
}
-// Const array of foo -> mapped 'to'
-
+// Const array of a const-qualified struct type.
+// Mapped as TO|TARGET_PARAM|IMPLICIT = 545.
// LABEL: define dso_local void @_Z16test_const_arrayv
// CHECK: store ptr @.offload_maptypes.10, ptr {{.*}}, align 8
void test_const_array() {
@@ -111,11 +122,14 @@ void test_const_array() {
}
}
-// Explicit map(tofrom:) on a const struct -> mapped 'tofrom'
+// ---------------------------------------------------------------------------
+// Explicit map clause tests
+// ---------------------------------------------------------------------------
-// LABEL: define dso_local void @_Z27test_explicit_map_overridesv
+// Explicit map(tofrom:) on a const struct. Mapped as TO|TARGET_PARAM = 33.
+// LABEL: define dso_local void @_Z27test_explicit_tofrom_const
// CHECK: store ptr @.offload_maptypes.12, ptr {{.*}}, align 8
-void test_explicit_map_overrides() {
+void test_explicit_tofrom_const() {
const foo a(2);
#pragma omp target map(tofrom:a)
{
@@ -123,10 +137,69 @@ void test_explicit_map_overrides() {
}
}
-// Mixed: const foo (to) and non-const foo (tofrom) in the same region.
+// Explicit map(from:) on a const struct. The FROM clause is ignored.
+// Mapped as TARGET_PARAM = 32.
+// LABEL: define dso_local void @_Z24test_explicit_from_constv
+// CHECK: store ptr @.offload_maptypes.14, ptr {{.*}}, align 8
+void test_explicit_from_const() {
+ const foo a(2);
+#pragma omp target map(from:a)
+ {
+ int x = a.i;
+ }
+}
+
+// Explicit map(to:) on a const struct. Mapped as TO|TARGET_PARAM = 33.
+// LABEL: define dso_local void @_Z22test_explicit_to_constv()
+// CHECK: store ptr @.offload_maptypes.16, ptr {{.*}}, align 8
+void test_explicit_to_const() {
+ const foo a(2);
+#pragma omp target map(to:a)
+ {
+ int x = a.i;
+ }
+}
+
+// ---------------------------------------------------------------------------
+// target update from tests
+// ---------------------------------------------------------------------------
+
+// target update from on a const struct with no mutable members. The FROM clause
+// is ignored. Mapped as FROM = 2.
+// LABEL: define dso_local void @_Z29test_target_update_from_constv
+// CHECK: call void @__tgt_target_data_update_mapper(ptr @1, i64 -1, i32 1, ptr %3, ptr %4, ptr @.offload_sizes.17, ptr @.offload_maptypes.18, ptr null, ptr null)
+void test_target_update_from_const() {
+ const foo a(2);
+#pragma omp target update from(a)
+}
+
+// target update from on a non-const struct. Mapped as FROM = 2.
+// LABEL: define dso_local void @_Z32test_target_update_from_nonconstv
+// CHECK: call void @__tgt_target_data_update_mapper(ptr @1, i64 -1, i32 1, ptr %3, ptr %4, ptr @.offload_sizes.19, ptr @.offload_maptypes.20, ptr null, ptr null)
+void test_target_update_from_nonconst() {
+ foo a(2);
+#pragma omp target update from(a)
+}
+
+// target update from on a const struct that has a mutable member. Mapped as FROM = 2.
+// LABEL: define dso_local void @_Z37test_target_update_from_const_mutablev
+// CHECK: call void @__tgt_target_data_update_mapper(ptr @1, i64 -1, i32 1, ptr %3, ptr %4, ptr @.offload_sizes.21, ptr @.offload_maptypes.22, ptr null, ptr null)
+
+void test_target_update_from_const_mutable() {
+ const foo_mutable a(2);
+#pragma omp target update from(a)
+}
+// ---------------------------------------------------------------------------
+// Combined tests
+// ---------------------------------------------------------------------------
+
+// Mixed region with one const and one non-const variable of the same struct
+// type. Each variable gets its own map type: const maps as
+// TO|TARGET_PARAM|IMPLICIT = 545, non-const maps as
+// TO|FROM|TARGET_PARAM|IMPLICIT = 547.
// LABEL: define dso_local void @_Z10test_mixedv
-// CHECK: store ptr @.offload_maptypes.14, ptr {{.*}}, align 8
+// CHECK: store ptr @.offload_maptypes.24, ptr {{.*}}, align 8
void test_mixed() {
const foo ca(2);
foo ma(3);
@@ -137,10 +210,10 @@ void test_mixed() {
}
}
-// Defaultmap(tofrom:aggregate) explicit -> mapped 'to'.
-
+// Explicit defaultmap(tofrom:aggregate) directive on a const struct.
+// Mapped as TO|TARGET_PARAM|IMPLICIT = 545.
// LABEL: define dso_local void @_Z31test_defaultmap_tofrom_explicitv
-// CHECK: store ptr @.offload_maptypes.16, ptr {{.*}}, align 8
+// CHECK: store ptr @.offload_maptypes.26, ptr {{.*}}, align 8
void test_defaultmap_tofrom_explicit() {
const foo a(2);
#pragma omp target defaultmap(tofrom:aggregate)
More information about the cfe-commits
mailing list