9e62802422
This patch implements some parts of the target variable mapping changes specified in OpenMP 5.0, including base-pointer attachment/detachment behavior for array section list-items in map clauses, and ordering of map clauses according to map kind. 2020-11-10 Chung-Lin Tang <cltang@codesourcery.com> gcc/c-family/ChangeLog: * c-common.h (c_omp_adjust_map_clauses): New declaration. * c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses. (c_omp_adjust_map_clauses): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/ChangeLog: * gimplify.c (is_or_contains_p): New static helper function. (omp_target_reorder_clauses): New function. (gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to reorder clause list according to OpenMP 5.0 rules. Add handling of GOMP_MAP_ATTACH_DETACH for OpenMP cases. * omp-low.c (is_omp_target): New static helper function. (scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. (lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. gcc/testsuite/ChangeLog: * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid. * gfortran.dg/gomp/map-2.f90: Likewise. * c-c++-common/gomp/map-5.c: New testcase. libgomp/ChangeLog: * libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag usable. * oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to 'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'. (goacc_enter_datum): Likewise for call to gomp_map_vars_async. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_internal): Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use of gomp_attach_pointer for OpenMP cases. (gomp_exit_data): Add handling of GOMP_MAP_DETACH. (GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH. * testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.
82 lines
1.2 KiB
C
82 lines
1.2 KiB
C
#include <stdlib.h>
|
|
|
|
struct S
|
|
{
|
|
int a, b;
|
|
int *ptr;
|
|
int c, d;
|
|
};
|
|
typedef struct S S;
|
|
|
|
#pragma omp declare target
|
|
int *gp;
|
|
#pragma omp end declare target
|
|
|
|
#define N 10
|
|
int main (void)
|
|
{
|
|
/* Test to see if pointer attachment works, for scalar pointers,
|
|
and pointer fields in structures. */
|
|
|
|
int *ptr = (int *) malloc (sizeof (int) * N);
|
|
int *orig_ptr = ptr;
|
|
|
|
#pragma omp target map (ptr, ptr[:N])
|
|
{
|
|
for (int i = 0; i < N; i++)
|
|
ptr[i] = N - i;
|
|
}
|
|
|
|
if (ptr != orig_ptr)
|
|
abort ();
|
|
|
|
for (int i = 0; i < N; i++)
|
|
if (ptr[i] != N - i)
|
|
abort ();
|
|
|
|
S s = { 0 };
|
|
s.ptr = ptr;
|
|
#pragma omp target map (s, s.ptr[:N])
|
|
{
|
|
for (int i = 0; i < N; i++)
|
|
s.ptr[i] = i;
|
|
|
|
s.a = 1;
|
|
s.b = 2;
|
|
}
|
|
|
|
if (s.ptr != ptr)
|
|
abort ();
|
|
|
|
for (int i = 0; i < N; i++)
|
|
if (s.ptr[i] != i)
|
|
abort ();
|
|
|
|
if (s.a != 1 || s.b != 2 || s.c != 0 || s.d != 0)
|
|
abort ();
|
|
|
|
gp = (int *) malloc (sizeof (int) * N);
|
|
orig_ptr = gp;
|
|
|
|
for (int i = 0; i < N; i++)
|
|
gp[i] = i - 1;
|
|
|
|
#pragma omp target map (gp[:N])
|
|
{
|
|
for (int i = 0; i < N; i++)
|
|
gp[i] += 1;
|
|
}
|
|
|
|
if (gp != orig_ptr)
|
|
abort ();
|
|
|
|
for (int i = 0; i < N; i++)
|
|
if (gp[i] != i)
|
|
abort ();
|
|
|
|
free (ptr);
|
|
free (gp);
|
|
|
|
return 0;
|
|
}
|
|
|