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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
