[OpenMP] Map const-qualified target map variables as 'to'. (#185918)
This patch updates the mapping kind for const-qualified variables from`tofrom` to `to`, ensuring correct and standards-compliant mapping semantics for const variables.
This commit is contained in:
Zahira Ammarguellat
GitHub
parent
c949c9be61
commit
cf89c33e6d
@ -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,36 @@ 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())
|
||||
// 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;
|
||||
}
|
||||
|
||||
namespace {
|
||||
struct VariableImplicitInfo {
|
||||
static const unsigned MapKindNum = OMPC_MAP_unknown;
|
||||
@ -4128,7 +4162,8 @@ public:
|
||||
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 +4260,8 @@ public:
|
||||
OpenMPDefaultmapClauseKind ClauseKind =
|
||||
getVariableCategoryFromDecl(SemaRef.getLangOpts(), FD);
|
||||
OpenMPMapClauseKind Kind = getMapClauseKindFromModifier(
|
||||
Modifier, /*IsAggregateOrDeclareTarget=*/true);
|
||||
Modifier, /*IsAggregateOrDeclareTarget=*/true,
|
||||
/*HasConstQualifier=*/false);
|
||||
ImpInfo.Mappings[ClauseKind][Kind].insert(E);
|
||||
return;
|
||||
}
|
||||
@ -23498,6 +23534,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,
|
||||
|
||||
246
clang/test/OpenMP/map_const_aggregate.cpp
Normal file
246
clang/test/OpenMP/map_const_aggregate.cpp
Normal file
@ -0,0 +1,246 @@
|
||||
// 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
|
||||
|
||||
// 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;
|
||||
};
|
||||
|
||||
// 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;
|
||||
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 [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 [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]
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// 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() {
|
||||
const foo a(2);
|
||||
#pragma omp target
|
||||
{
|
||||
int x = a.i;
|
||||
}
|
||||
}
|
||||
|
||||
// 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() {
|
||||
foo a(2);
|
||||
#pragma omp target
|
||||
{
|
||||
int x = a.i;
|
||||
}
|
||||
}
|
||||
|
||||
// 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() {
|
||||
const foo_mutable a(2);
|
||||
#pragma omp target
|
||||
{
|
||||
a.m = 1;
|
||||
}
|
||||
}
|
||||
|
||||
// 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() {
|
||||
const foo_nested a(2);
|
||||
#pragma omp target
|
||||
{
|
||||
int x = a.inner.i;
|
||||
}
|
||||
}
|
||||
|
||||
// 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() {
|
||||
const foo_nested_mutable a(2);
|
||||
#pragma omp target
|
||||
{
|
||||
a.inner.m = 1;
|
||||
}
|
||||
}
|
||||
|
||||
// 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() {
|
||||
const foo arr[4] = {1, 2, 3, 4};
|
||||
#pragma omp target
|
||||
{
|
||||
int x = arr[0].i;
|
||||
}
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Explicit map clause tests
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
// 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_tofrom_const() {
|
||||
const foo a(2);
|
||||
#pragma omp target map(tofrom:a)
|
||||
{
|
||||
int x = a.i;
|
||||
}
|
||||
}
|
||||
|
||||
// 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.24, ptr {{.*}}, align 8
|
||||
void test_mixed() {
|
||||
const foo ca(2);
|
||||
foo ma(3);
|
||||
#pragma omp target
|
||||
{
|
||||
int x = ca.i;
|
||||
ma.i = 99;
|
||||
}
|
||||
}
|
||||
|
||||
// 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.26, ptr {{.*}}, align 8
|
||||
void test_defaultmap_tofrom_explicit() {
|
||||
const foo a(2);
|
||||
#pragma omp target defaultmap(tofrom:aggregate)
|
||||
{
|
||||
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)
|
||||
}
|
||||
Loading…
x
Reference in New Issue
Block a user