38944ec2a6
Reduce number of enum values passed to libgomp as
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} have the same semantic as
GOMP_MAP_FORCE_PRESENT (i.e. abort if not present, otherwise ignore);
that's different to GOMP_MAP_ALWAYS_PRESENT_{TO,TOFROM,FROM} which also
abort if not present but copy data when present. This is is a follow-up to
the commit r14-1579-g4ede915d5dde93 done 6 days ago.
Additionally, the commit improves a libgomp run-time and a C/C++ compile-time
error wording and extends testcases a tiny bit.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_map): Reword error message for
clearness especially with 'omp target (enter/exit) data.'
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_map): Reword error message for
clearness especially with 'omp target (enter/exit) data.'
* semantics.cc (handle_omp_array_sections): Handle
GOMP_MAP_{ALWAYS_,}PRESENT_{TO,TOFROM,FROM,ALLOC} enum values.
gcc/ChangeLog:
* gimplify.cc (gimplify_adjust_omp_clauses_1): Use
GOMP_MAP_FORCE_PRESENT for 'present alloc' implicit mapping.
(gimplify_adjust_omp_clauses): Change
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to the equivalent
GOMP_MAP_FORCE_PRESENT.
* omp-low.cc (lower_omp_target): Remove handling of no-longer valid
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC}; update map kinds used for
to/from clauses with present modifier.
include/ChangeLog:
* gomp-constants.h (enum gomp_map_kind): Change the enum values
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to be compiler only.
(GOMP_MAP_PRESENT_P): Update to include also GOMP_MAP_FORCE_PRESENT.
libgomp/ChangeLog:
* target.c (gomp_to_device_kind_p, gomp_map_vars_internal): Replace
GOMP_MAP_PRESENT_{FROM,TO,TOFROM,ACLLOC} by GOMP_MAP_FORCE_PRESENT.
(gomp_map_vars_internal, gomp_update): Likewise; unify and improve
error message.
* testsuite/libgomp.c-c++-common/target-present-2.c: Update for
changed error message.
* testsuite/libgomp.fortran/target-present-1.f90: Likewise.
* testsuite/libgomp.fortran/target-present-2.f90: Likewise.
* testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.
* testsuite/libgomp.c-c++-common/target-present-1.c: Likewise and
extend testcase to check that data is copied when needed.
* testsuite/libgomp.c-c++-common/target-present-3.c: Likewise.
* testsuite/libgomp.fortran/target-present-3.f90: Likewise.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/defaultmap-4.c: Update scan-tree-dump.
* c-c++-common/gomp/map-9.c: Likewise.
* gfortran.dg/gomp/defaultmap-8.f90: Likewise.
* gfortran.dg/gomp/map-11.f90: Likewise.
* gfortran.dg/gomp/target-update-1.f90: Likewise.
* gfortran.dg/gomp/map-12.f90: Likewise; also check original dump.
* c-c++-common/gomp/map-6.c: Update dg-error and also check
clause error with 'target (enter/exit) data'.
40 lines
1.2 KiB
C
40 lines
1.2 KiB
C
#include <stdio.h>
|
|
|
|
#define N 100
|
|
|
|
int main (void)
|
|
{
|
|
int a[N], b[N], c[N], d[N];
|
|
|
|
for (int i = 0; i < N; i++) {
|
|
a[i] = i * 2;
|
|
b[i] = i * 3 + 1;
|
|
d[i] = i * 5;
|
|
}
|
|
|
|
#pragma omp target enter data map (alloc: c, d) map(to: a)
|
|
#pragma omp target map (present, always, to: d)
|
|
for (int i = 0; i < N; i++)
|
|
if (d[i] != i * 5)
|
|
__builtin_abort ();
|
|
|
|
/* a has already been mapped and 'c' allocated so this should be okay. */
|
|
#pragma omp target map (present, to: a) map(present, always, from: c)
|
|
for (int i = 0; i < N; i++)
|
|
c[i] = a[i];
|
|
|
|
for (int i = 0; i < N; i++)
|
|
if (c[i] != i * 2)
|
|
__builtin_abort ();
|
|
|
|
fprintf (stderr, "CheCKpOInT\n");
|
|
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
|
|
|
/* b has not been allocated, so this should result in an error. */
|
|
/* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
|
|
/* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */
|
|
#pragma omp target map (present, to: b)
|
|
for (int i = 0; i < N; i++)
|
|
c[i] += b[i];
|
|
#pragma omp target exit data map (from: c)
|
|
}
|