[clang] [OpenMP] Map const-qualified target map variables as 'to'. (PR #185918)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 11 09:48:16 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Zahira Ammarguellat (zahiraam)
<details>
<summary>Changes</summary>
This patch updates the mapping kind for const-qualified variables from`tofrom` to `to`, ensuring correct and standards-compliant mapping semantics for const variables.
---
Full diff: https://github.com/llvm/llvm-project/pull/185918.diff
2 Files Affected:
- (modified) clang/lib/Sema/SemaOpenMP.cpp (+34-4)
- (added) clang/test/OpenMP/map_const_aggregate.cpp (+150)
``````````diff
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index e8a7c1d5e9288..4a1003a406b35 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;
+ }
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/185918
More information about the cfe-commits
mailing list