https://github.com/zahiraam updated https://github.com/llvm/llvm-project/pull/185918
>From 310f7226b1dbc101823253d99103b4a300f73635 Mon Sep 17 00:00:00 2001 From: Ammarguellat <[email protected]> Date: Wed, 11 Mar 2026 09:38:48 -0700 Subject: [PATCH 1/5] [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 <[email protected]> Date: Wed, 11 Mar 2026 09:47:01 -0700 Subject: [PATCH 2/5] 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 <[email protected]> Date: Thu, 12 Mar 2026 07:21:14 -0700 Subject: [PATCH 3/5] 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) >From 56c17e96278cfe4de9d5c90b44a8cad18123a823 Mon Sep 17 00:00:00 2001 From: Ammarguellat <[email protected]> Date: Tue, 17 Mar 2026 06:03:27 -0700 Subject: [PATCH 4/5] Addressed review comments --- clang/lib/Sema/SemaOpenMP.cpp | 6 ++++++ clang/test/OpenMP/map_const_aggregate.cpp | 25 ++++++++++++++++++++++- 2 files changed, 30 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 05b6e31751a96..2d6a952e07ca4 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -3930,6 +3930,12 @@ static bool hasConstQualifiedMappingType(QualType T) { if (!T.isConstQualified()) return false; if (const auto *RD = T->getAsCXXRecordDecl()) + // TODO : Per OpenMP 6.0 p299 lines 3-4, non-mutable members of a + // const-qualified struct should also be ignored for 'from'. This + // requires per-member mapping granularity via compiler-generated + // default mappers and a mechanism to ensure constness to the mapper. + // For now we conservatively treat any struct with mutable members as + // requiring full 'tofrom'. return hasNoMutableFields(RD); return true; } diff --git a/clang/test/OpenMP/map_const_aggregate.cpp b/clang/test/OpenMP/map_const_aggregate.cpp index ed9d67bc4c064..9a8f3e85b7ac9 100644 --- a/clang/test/OpenMP/map_const_aggregate.cpp +++ b/clang/test/OpenMP/map_const_aggregate.cpp @@ -22,6 +22,10 @@ struct foo_mutable { mutable int m; }; +// TODO: A const foo_mutable should ideally only copy back its mutable +// member 'm' and ignore non-mutable member 'i' on a 'from' mapping, per +// OpenMP 6.0 p299 lines 3-4. This requires per-member mapper generation +// and is left for a follow-up patch. struct foo_nested { foo_nested(int j) : inner(j), z(j) {}; foo inner; @@ -48,6 +52,7 @@ struct foo_nested_mutable { // 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] +// CHECK: @.offload_maptypes.28 = private unnamed_addr constant [1 x i64] [i64 2] // --------------------------------------------------------------------------- // Implicit mapping tests (no explicit map clause, defaultmap governs) @@ -184,7 +189,6 @@ void test_target_update_from_nonconst() { // 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) @@ -221,3 +225,22 @@ void test_defaultmap_tofrom_explicit() { int x = a.i; } } + +// User-defined mapper on const struct — FROM must NOT be suppressed because the +// mapper accesses non-const pointee data py[0:10]. +// Mapped as FROM = 2. +// LABEL: define dso_local void @_Z30test_user_defined_mapper_constv +// CHECK: call void @__tgt_target_data_update_mapper(ptr @1, i64 -1, i32 1, ptr {{.*}}, ptr {{.*}}, ptr @.offload_sizes.27, ptr @.offload_maptypes.28, ptr null, ptr {{.*}}) +int y[10]; +struct S { + int x; + int *py; +}; + +#pragma omp declare mapper(m1: const S s) map(alloc: s.x, s.py) map(from: s.py[0:10]) + +void test_user_defined_mapper_const() { + int data[10] = {0}; + const S s1 = {1, data}; + #pragma omp target update from(mapper(m1): s1) +} >From 5fe58c502c8349c1b560c311931dc0d91ac556a4 Mon Sep 17 00:00:00 2001 From: Ammarguellat <[email protected]> Date: Tue, 17 Mar 2026 10:02:57 -0700 Subject: [PATCH 5/5] Sync the LIT test --- clang/test/OpenMP/map_const_aggregate.cpp | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/clang/test/OpenMP/map_const_aggregate.cpp b/clang/test/OpenMP/map_const_aggregate.cpp index 9a8f3e85b7ac9..fbb1c3a5abe60 100644 --- a/clang/test/OpenMP/map_const_aggregate.cpp +++ b/clang/test/OpenMP/map_const_aggregate.cpp @@ -38,20 +38,20 @@ struct foo_nested_mutable { 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 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 = private unnamed_addr constant [2 x i64] [i64 545, i64 288] +// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [2 x i64] [i64 547, i64 288] +// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 547, i64 288] +// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 545, i64 288] +// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [2 x i64] [i64 547, i64 288] +// CHECK: @.offload_maptypes.10 = private unnamed_addr constant [2 x i64] [i64 545, i64 288] +// CHECK: @.offload_maptypes.12 = private unnamed_addr constant [2 x i64] [i64 33, i64 288] +// CHECK: @.offload_maptypes.14 = private unnamed_addr constant [2 x i64] [i64 32, i64 288] +// CHECK: @.offload_maptypes.16 = private unnamed_addr constant [2 x i64] [i64 33, i64 288] // 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] +// CHECK: @.offload_maptypes.24 = private unnamed_addr constant [3 x i64] [i64 545, i64 547, i64 288] +// CHECK: @.offload_maptypes.26 = private unnamed_addr constant [2 x i64] [i64 545, i64 288] // CHECK: @.offload_maptypes.28 = private unnamed_addr constant [1 x i64] [i64 2] // --------------------------------------------------------------------------- _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
