6c0399378e
This patch implements three pieces of functionality:
(1) Adjust array section mapping to have standards conforming behavior,
mapping array sections should *NOT* also map the base-pointer:
struct S { int *ptr; ... };
struct S s;
Instead of generating this during gimplify:
map(to:*_1 [len: 400]) map(attach:s.ptr [bias: 0])
Now, adjust to:
(i.e. do not map the base-pointer together. The attach operation is still
generated, and if s.ptr is already mapped prior, attachment will happen)
The correct way of achieving the base-pointer-also-mapped behavior would be to
use:
(A small Fortran front-end patch to trans-openmp.c:gfc_trans_omp_array_section
is also included, which removes generation of a GOMP_MAP_ALWAYS_POINTER for
array types, which appears incorrect and causes a regression in
libgomp.fortranlibgomp.fortran/struct-elem-map-1.f90)
(2) Related to the first item above, are fixes in libgomp/target.c to not
overwrite attached pointers when handling device<->host copies, mainly for the
"always" case.
(3) The third is a set of changes to the C/C++ front-ends to extend the allowed
component access syntax in map clauses. These changes are enabled for both
OpenACC and OpenMP.
gcc/c/ChangeLog:
* c-parser.c (struct omp_dim): New struct type for use inside
c_parser_omp_variable_list.
(c_parser_omp_variable_list): Allow multiple levels of array and
component accesses in array section base-pointer expression.
(c_parser_omp_clause_to): Set 'allow_deref' to true in call to
c_parser_omp_var_list_parens.
(c_parser_omp_clause_from): Likewise.
* c-typeck.c (handle_omp_array_sections_1): Extend allowed range
of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
POINTER_PLUS_EXPR.
(c_finish_omp_clauses): Extend allowed ranged of expressions
involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.
gcc/cp/ChangeLog:
* parser.c (struct omp_dim): New struct type for use inside
cp_parser_omp_var_list_no_open.
(cp_parser_omp_var_list_no_open): Allow multiple levels of array and
component accesses in array section base-pointer expression.
(cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to
cp_parser_omp_var_list for to/from clauses.
* semantics.c (handle_omp_array_sections_1): Extend allowed range
of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
POINTER_PLUS_EXPR.
(handle_omp_array_sections): Adjust pointer map generation of
references.
(finish_omp_clauses): Extend allowed ranged of expressions
involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.
gcc/fortran/ChangeLog:
* trans-openmp.c (gfc_trans_omp_array_section): Do not generate
GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type.
gcc/ChangeLog:
* gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter,
accomodate case where 'offset' return of get_inner_reference is
non-NULL.
(is_or_contains_p): Further robustify conditions.
(omp_target_reorder_clauses): In alloc/to/from sorting phase, also
move following GOMP_MAP_ALWAYS_POINTER maps along. Add new sorting
phase where we make sure pointers with an attach/detach map are ordered
correctly.
(gimplify_scan_omp_clauses): Add modifications to avoid creating
GOMP_MAP_STRUCT and associated alloc map for attach/detach maps.
gcc/testsuite/ChangeLog:
* c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase.
* c-c++-common/gomp/target-enter-data-1.c: New testcase.
* c-c++-common/gomp/target-implicit-map-2.c: New testcase.
libgomp/ChangeLog:
* target.c (gomp_map_vars_existing): Make sure attached pointer is
not overwritten during cross-host/device copying.
(gomp_update): Likewise.
(gomp_exit_data): Likewise.
* testsuite/libgomp.c++/target-11.C: Adjust testcase.
* testsuite/libgomp.c++/target-12.C: Likewise.
* testsuite/libgomp.c++/target-15.C: Likewise.
* testsuite/libgomp.c++/target-16.C: Likewise.
* testsuite/libgomp.c++/target-17.C: Likewise.
* testsuite/libgomp.c++/target-21.C: Likewise.
* testsuite/libgomp.c++/target-23.C: Likewise.
* testsuite/libgomp.c/target-23.c: Likewise.
* testsuite/libgomp.c/target-29.c: Likewise.
* testsuite/libgomp.c-c++-common/target-implicit-map-2.c: New testcase.
174 lines
2.9 KiB
C
174 lines
2.9 KiB
C
extern "C" void abort ();
|
|
struct T { char t[270]; };
|
|
struct S { int (&x)[10]; int *&y; T t; int &z; S (); ~S (); };
|
|
|
|
template <int N>
|
|
void
|
|
foo (S s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x[0:N], s.y, s.y[0:N]) map (s.t.t[16:3]) map (from: err)
|
|
{
|
|
err = s.x[2] != 28 || s.y[2] != 37 || s.t.t[17] != 81;
|
|
s.x[2]++;
|
|
s.y[2]++;
|
|
s.t.t[17]++;
|
|
}
|
|
if (err || s.x[2] != 29 || s.y[2] != 38 || s.t.t[17] != 82)
|
|
abort ();
|
|
}
|
|
|
|
template <int N>
|
|
void
|
|
bar (S s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x, s.z)map(from:err)
|
|
{
|
|
err = s.x[2] != 29 || s.z != 6;
|
|
s.x[2]++;
|
|
s.z++;
|
|
}
|
|
if (err || s.x[2] != 30 || s.z != 7)
|
|
abort ();
|
|
}
|
|
|
|
template <int N>
|
|
void
|
|
foo2 (S &s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x[N:10], s.y, s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
|
|
{
|
|
err = s.x[2] != 30 || s.y[2] != 38 || s.t.t[17] != 81;
|
|
s.x[2]++;
|
|
s.y[2]++;
|
|
s.t.t[17]++;
|
|
}
|
|
if (err || s.x[2] != 31 || s.y[2] != 39 || s.t.t[17] != 82)
|
|
abort ();
|
|
}
|
|
|
|
template <int N>
|
|
void
|
|
bar2 (S &s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x, s.z)map(from:err)
|
|
{
|
|
err = s.x[2] != 31 || s.z != 7;
|
|
s.x[2]++;
|
|
s.z++;
|
|
}
|
|
if (err || s.x[2] != 32 || s.z != 8)
|
|
abort ();
|
|
}
|
|
|
|
template <typename U>
|
|
void
|
|
foo3 (U s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
|
|
{
|
|
err = s.x[2] != 32 || s.y[2] != 39 || s.t.t[17] != 82;
|
|
s.x[2]++;
|
|
s.y[2]++;
|
|
s.t.t[17]++;
|
|
}
|
|
if (err || s.x[2] != 33 || s.y[2] != 40 || s.t.t[17] != 83)
|
|
abort ();
|
|
}
|
|
|
|
template <typename U>
|
|
void
|
|
bar3 (U s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x, s.z)map(from:err)
|
|
{
|
|
err = s.x[2] != 33 || s.z != 8;
|
|
s.x[2]++;
|
|
s.z++;
|
|
}
|
|
if (err || s.x[2] != 34 || s.z != 9)
|
|
abort ();
|
|
}
|
|
|
|
template <typename U>
|
|
void
|
|
foo4 (U &s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
|
|
{
|
|
err = s.x[2] != 34 || s.y[2] != 40 || s.t.t[17] != 82;
|
|
s.x[2]++;
|
|
s.y[2]++;
|
|
s.t.t[17]++;
|
|
}
|
|
if (err || s.x[2] != 35 || s.y[2] != 41 || s.t.t[17] != 83)
|
|
abort ();
|
|
}
|
|
|
|
template <typename U>
|
|
void
|
|
bar4 (U &s)
|
|
{
|
|
int err;
|
|
#pragma omp target map (s.x, s.z)map(from:err)
|
|
{
|
|
err = s.x[2] != 35 || s.z != 9;
|
|
s.x[2]++;
|
|
s.z++;
|
|
}
|
|
if (err || s.x[2] != 36 || s.z != 10)
|
|
abort ();
|
|
}
|
|
|
|
int xt[10] = { 1, 2, 28, 3, 4, 5, 6, 7, 8, 9 };
|
|
int yt[10] = { 1, 2, 37, 3, 4, 5, 6, 7, 8, 9 };
|
|
int *yp = yt;
|
|
int zt = 6;
|
|
|
|
S::S () : x (xt), y (yp), z (zt)
|
|
{
|
|
}
|
|
|
|
S::~S ()
|
|
{
|
|
}
|
|
|
|
int
|
|
main ()
|
|
{
|
|
S s;
|
|
s.t.t[16] = 5;
|
|
s.t.t[17] = 81;
|
|
s.t.t[18] = 9;
|
|
foo <10> (s);
|
|
if (s.t.t[17] != 81)
|
|
abort ();
|
|
bar <7> (s);
|
|
foo2 <0> (s);
|
|
if (s.t.t[17] != 82)
|
|
abort ();
|
|
bar2 <21> (s);
|
|
foo3 <S> (s);
|
|
if (s.t.t[17] != 82)
|
|
abort ();
|
|
bar3 <S> (s);
|
|
foo4 <S> (s);
|
|
if (s.t.t[17] != 83)
|
|
abort ();
|
|
bar4 <S> (s);
|
|
s.x[2] -= 4;
|
|
s.y[2] -= 2;
|
|
s.z -= 2;
|
|
s.t.t[17]--;
|
|
foo3 <S &> (s);
|
|
if (s.t.t[17] != 83)
|
|
abort ();
|
|
bar3 <S &> (s);
|
|
}
|