This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[patch, gomp4, committed] Adjust copy/copyin/copyout/create for OpenACC 2.5
- From: Chung-Lin Tang <chunglin_tang at mentor dot com>
- To: gcc-patches <gcc-patches at gcc dot gnu dot org>
- Cc: Thomas Schwinge <thomas at codesourcery dot com>, Cesar Philippidis <cesar at codesourcery dot com>
- Date: Thu, 13 Apr 2017 19:11:16 +0800
- Subject: [patch, gomp4, committed] Adjust copy/copyin/copyout/create for OpenACC 2.5
- Authentication-results: sourceware.org; auth=none
The behavior of the copy/copyin/copyout/create clauses has been changed
in OpenACC 2.5 to be like the present_or_* variants, and the original
present_or_* syntax relegated to legacy status.
This patch removes the presence of any PRESENT_OR_* symbols, and
changes the mapping of the copy/copyin/copyout/create clauses to map
kinds without the FORCE flag. Library routines acc_[present_or_]copy, etc.
has also been updated.
This patch has been applied to the gomp-4_0-branch.
Chung-Lin
2017-04-13 Chung-Lin Tang <cltang@codesourcery.com>
gcc/
* gimplify.c (gimplify_oacc_declare_1): Remove GOMP_MAP_FORCE_* cases.
gcc/c-family/
* c-pragma.h (enum pragma_omp_clauses):
Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Remove occurences of
PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE, adjust
them to non-PRESENT_OR values.
(c_parser_oacc_data_clause): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*
cases, remove FORCE from PRAGMA_OACC_CLAUSE_COPY/COPYIN/COPYOUT/CREATE
kinds, update comment description.
(c_parser_oacc_all_clauses): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*
cases.
(OACC_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*.
(OACC_DECLARE_CLAUSE_MASK): Likewise.
(c_parser_oacc_declare): Remove GOMP_MAP_FORCE_ALLOC/TO, change to
COMP_MAP_ALLOC/TO.
(OACC_ENTER_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*.
(OACC_KERNELS_CLAUSE_MASK): Likewise.
(OACC_PARALLEL_CLAUSE_MASK): Likewise.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Remove occurences of
PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE, adjust
them to non-PRESENT_OR values.
(cp_parser_oacc_data_clause): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*
cases, remove FORCE from PRAGMA_OACC_CLAUSE_COPY/COPYIN/COPYOUT/CREATE
kinds, update comment description.
(cp_parser_oacc_all_clauses): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*
cases.
(OACC_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*.
(OACC_DECLARE_CLAUSE_MASK): Likewise.
(cp_parser_oacc_declare): Remove GOMP_MAP_FORCE_ALLOC/TO, change to
COMP_MAP_ALLOC/TO.
(OACC_ENTER_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*.
(OACC_KERNELS_CLAUSE_MASK): Likewise.
(OACC_PARALLEL_CLAUSE_MASK): Likewise.
gcc/fortran/
* openmp.c (enum omp_mask2):
Remove OMP_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE.
(gfc_match_omp_clauses): Remove FORCE flag from OpenACC OMP_MAP_* cases,
adjust and remove PRESENT_OR_ values.
(OACC_PARALLEL_CLAUSES): Remove OMP_CLAUSE_PRESENT_OR_*.
(OACC_KERNELS_CLAUSES): Likewise.
(OACC_DATA_CLAUSES): Likewise.
(OACC_DECLARE_CLAUSES): Likewise.
(OACC_ENTER_DATA_CLAUSES): Likewise.
gcc/testsuite/
* gfortran.dg/goacc/data-tree.f95: Remove "force_" from dump scan.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/reduction-promotions.f90: Likewise.
* gfortran.dg/goacc/combined-directives.f90: Likewise.
* gfortran.dg/goacc/default-4.f: Likewise.
* gfortran.dg/declare-2.f95: Adjust error scan.
* c-c++-common/goacc/acc-data-chain.c: Remove "force_" from dump scan.
* c-c++-common/goacc/default-4.c: Likewise.
* c-c++-common/goacc/declare-2.c: Move present_or_copyin/create tests
from here to...
* c-c++-common/goacc/declare-1.c: ...here.
libgomp/
* oacc-mem.c (acc_create): Add FLAG_PRESENT to call of
present_create_copy.
(acc_create_async): Likewise.
(acc_copyin): Likewise.
(acc_copyin_async): Likewise.
(acc_present_or_create): Remove definition and change to alias of
acc_create.
(acc_present_or_copyin): Remove definition and change to alias of
acc_copyin.
* oacc-parallel.c (GOACC_enter_exit_data): Add GOMP_MAP_FROM as
handled map kind.
* testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Delete.
* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Delete.
* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Delete.
* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Delete.
* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Delete.
* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Delete.
* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Delete.
* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Delete.
* testsuite/libgomp.oacc-fortran/data-already-1.f: Delete.
* testsuite/libgomp.oacc-fortran/data-already-2.f: Delete.
* testsuite/libgomp.oacc-fortran/data-already-3.f: Delete.
* testsuite/libgomp.oacc-fortran/data-already-4.f: Delete.
* testsuite/libgomp.oacc-fortran/data-already-5.f: Delete.
* testsuite/libgomp.oacc-fortran/data-already-6.f: Delete.
* testsuite/libgomp.oacc-fortran/data-already-7.f: Delete.
* testsuite/libgomp.oacc-fortran/data-already-8.f: Delete.
* testsuite/libgomp.oacc-c-c++-common/lib-16.c: Remove dg-shouldfail.
* testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 246808)
+++ gcc/gimplify.c (working copy)
@@ -8584,24 +8584,10 @@ gimplify_oacc_declare_1 (tree clause)
switch (kind)
{
case GOMP_MAP_ALLOC:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
new_op = GOMP_MAP_DELETE;
ret = true;
break;
- case GOMP_MAP_FORCE_FROM:
- OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
- new_op = GOMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_FORCE_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO);
- new_op = GOMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
case GOMP_MAP_FROM:
OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
new_op = GOMP_MAP_FROM;
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c (revision 246808)
+++ gcc/c/c-parser.c (working copy)
@@ -10447,18 +10447,20 @@ c_parser_omp_clause_name (c_parser *parser, bool c
result = PRAGMA_OMP_CLAUSE_PARALLEL;
else if (!strcmp ("present", p))
result = PRAGMA_OACC_CLAUSE_PRESENT;
+ /* As of OpenACC 2.5, these are now aliases of the non-present_or
+ clauses. */
else if (!strcmp ("present_or_copy", p)
|| !strcmp ("pcopy", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
+ result = PRAGMA_OACC_CLAUSE_COPY;
else if (!strcmp ("present_or_copyin", p)
|| !strcmp ("pcopyin", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
+ result = PRAGMA_OACC_CLAUSE_COPYIN;
else if (!strcmp ("present_or_copyout", p)
|| !strcmp ("pcopyout", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
+ result = PRAGMA_OACC_CLAUSE_COPYOUT;
else if (!strcmp ("present_or_create", p)
|| !strcmp ("pcreate", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
+ result = PRAGMA_OACC_CLAUSE_CREATE;
else if (!strcmp ("priority", p))
result = PRAGMA_OMP_CLAUSE_PRIORITY;
else if (!strcmp ("private", p))
@@ -10752,7 +10754,7 @@ c_parser_omp_var_list_parens (c_parser *parser, en
return list;
}
-/* OpenACC 2.0:
+/* OpenACC 2.5:
copy ( variable-list )
copyin ( variable-list )
copyout ( variable-list )
@@ -10760,15 +10762,7 @@ c_parser_omp_var_list_parens (c_parser *parser, en
delete ( variable-list )
device_resident ( variable-list )
link ( variable-list )
- present ( variable-list )
- present_or_copy ( variable-list )
- pcopy ( variable-list )
- present_or_copyin ( variable-list )
- pcopyin ( variable-list )
- present_or_copyout ( variable-list )
- pcopyout ( variable-list )
- present_or_create ( variable-list )
- pcreate ( variable-list ) */
+ present ( variable-list ) */
static tree
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
@@ -10778,16 +10772,16 @@ c_parser_oacc_data_clause (c_parser *parser, pragm
switch (c_kind)
{
case PRAGMA_OACC_CLAUSE_COPY:
- kind = GOMP_MAP_FORCE_TOFROM;
+ kind = GOMP_MAP_TOFROM;
break;
case PRAGMA_OACC_CLAUSE_COPYIN:
- kind = GOMP_MAP_FORCE_TO;
+ kind = GOMP_MAP_TO;
break;
case PRAGMA_OACC_CLAUSE_COPYOUT:
- kind = GOMP_MAP_FORCE_FROM;
+ kind = GOMP_MAP_FROM;
break;
case PRAGMA_OACC_CLAUSE_CREATE:
- kind = GOMP_MAP_FORCE_ALLOC;
+ kind = GOMP_MAP_ALLOC;
break;
case PRAGMA_OACC_CLAUSE_DELETE:
kind = GOMP_MAP_DELETE;
@@ -10807,18 +10801,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragm
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- kind = GOMP_MAP_TOFROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- kind = GOMP_MAP_TO;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- kind = GOMP_MAP_FROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- kind = GOMP_MAP_ALLOC;
- break;
default:
gcc_unreachable ();
}
@@ -13285,22 +13267,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_c
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "present";
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copy";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyin";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyout";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_create";
- break;
case PRAGMA_OACC_CLAUSE_PRIVATE:
clauses = c_parser_omp_clause_private (parser, clauses);
c_name = "private";
@@ -13706,11 +13672,7 @@ c_parser_oacc_cache (location_t loc, c_parser *par
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) )
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
static tree
c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p)
@@ -13741,11 +13703,7 @@ c_parser_oacc_data (location_t loc, c_parser *pars
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) )
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
static void
c_parser_oacc_declare (c_parser *parser)
@@ -13780,8 +13738,8 @@ c_parser_oacc_declare (c_parser *parser)
switch (OMP_CLAUSE_MAP_KIND (t))
{
case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
break;
@@ -13891,8 +13849,6 @@ c_parser_oacc_declare (c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
#define OACC_EXIT_DATA_CLAUSE_MASK \
@@ -14051,10 +14007,6 @@ c_parser_oacc_loop (location_t loc, c_parser *pars
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
#define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \
@@ -14076,10 +14028,6 @@ c_parser_oacc_loop (location_t loc, c_parser *pars
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
Index: gcc/c-family/c-pragma.h
===================================================================
--- gcc/c-family/c-pragma.h (revision 246808)
+++ gcc/c-family/c-pragma.h (working copy)
@@ -164,10 +164,6 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE,
PRAGMA_OACC_CLAUSE_SEQ,
PRAGMA_OACC_CLAUSE_TILE,
PRAGMA_OACC_CLAUSE_USE_DEVICE,
Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c (revision 246808)
+++ gcc/fortran/openmp.c (working copy)
@@ -818,10 +818,6 @@ enum omp_mask2
OMP_CLAUSE_COPYOUT,
OMP_CLAUSE_CREATE,
OMP_CLAUSE_PRESENT,
- OMP_CLAUSE_PRESENT_OR_COPY,
- OMP_CLAUSE_PRESENT_OR_COPYIN,
- OMP_CLAUSE_PRESENT_OR_COPYOUT,
- OMP_CLAUSE_PRESENT_OR_CREATE,
OMP_CLAUSE_DEVICEPTR,
OMP_CLAUSE_GANG,
OMP_CLAUSE_WORKER,
@@ -1082,7 +1078,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TOFROM, openacc,
+ OMP_MAP_TOFROM, openacc,
allow_derived))
continue;
if (mask & OMP_CLAUSE_COPYIN)
@@ -1091,7 +1087,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
{
if (gfc_match ("copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TO, true,
+ OMP_MAP_TO, true,
allow_derived))
continue;
}
@@ -1103,7 +1099,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM, true,
+ OMP_MAP_FROM, true,
allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYPRIVATE)
@@ -1114,7 +1110,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_ALLOC, true,
+ OMP_MAP_ALLOC, true,
allow_derived))
continue;
break;
@@ -1616,22 +1612,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
}
break;
case 'p':
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPY)
+ if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("pcopy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TOFROM, true, allow_derived))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN)
+ if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("pcopyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TO, true, allow_derived))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT)
+ if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("pcopyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_FROM, true, allow_derived))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE)
+ if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("pcreate ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_ALLOC, true, allow_derived))
@@ -1642,22 +1638,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
OMP_MAP_FORCE_PRESENT, false,
allow_derived))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPY)
+ if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("present_or_copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TOFROM, true, allow_derived))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN)
+ if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("present_or_copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TO, true, allow_derived))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT)
+ if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("present_or_copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_FROM, true, allow_derived))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE)
+ if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("present_or_create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_ALLOC, true, allow_derived))
@@ -2040,24 +2036,18 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \
- | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT \
- | OMP_CLAUSE_DEVICE_TYPE)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
+ | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
+ | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT \
- | OMP_CLAUSE_DEVICE_TYPE)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
+ | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
- | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE)
+ | OMP_CLAUSE_PRESENT)
#define OACC_LOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
| OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \
@@ -2071,16 +2061,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m
#define OACC_DECLARE_CLAUSES \
(omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT \
- | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_LINK)
+ | OMP_CLAUSE_PRESENT | OMP_CLAUSE_LINK)
#define OACC_UPDATE_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_SELF \
| OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_ENTER_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
- | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT_OR_COPYIN \
- | OMP_CLAUSE_PRESENT_OR_CREATE)
+ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
#define OACC_EXIT_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
| OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE)
@@ -2203,8 +2190,7 @@ gfc_match_oacc_declare (void)
if (s->ns->proc_name && s->ns->proc_name->attr.proc == PROC_MODULE)
{
- if (n->u.map_op != OMP_MAP_FORCE_ALLOC
- && n->u.map_op != OMP_MAP_FORCE_TO)
+ if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
{
gfc_error ("Invalid clause in module with $!ACC DECLARE at %L",
&where);
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c (revision 246808)
+++ gcc/cp/parser.c (working copy)
@@ -29887,18 +29887,20 @@ cp_parser_omp_clause_name (cp_parser *parser, bool
result = PRAGMA_OMP_CLAUSE_PARALLEL;
else if (!strcmp ("present", p))
result = PRAGMA_OACC_CLAUSE_PRESENT;
+ /* As of OpenACC 2.5, these are now aliases of the non-present_or
+ clauses. */
else if (!strcmp ("present_or_copy", p)
|| !strcmp ("pcopy", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
+ result = PRAGMA_OACC_CLAUSE_COPY;
else if (!strcmp ("present_or_copyin", p)
|| !strcmp ("pcopyin", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
+ result = PRAGMA_OACC_CLAUSE_COPYIN;
else if (!strcmp ("present_or_copyout", p)
|| !strcmp ("pcopyout", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
+ result = PRAGMA_OACC_CLAUSE_COPYOUT;
else if (!strcmp ("present_or_create", p)
|| !strcmp ("pcreate", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
+ result = PRAGMA_OACC_CLAUSE_CREATE;
else if (!strcmp ("priority", p))
result = PRAGMA_OMP_CLAUSE_PRIORITY;
else if (!strcmp ("proc_bind", p))
@@ -30169,7 +30171,7 @@ cp_parser_omp_var_list (cp_parser *parser, enum om
return list;
}
-/* OpenACC 2.0:
+/* OpenACC 2.5:
copy ( variable-list )
copyin ( variable-list )
copyout ( variable-list )
@@ -30178,15 +30180,7 @@ cp_parser_omp_var_list (cp_parser *parser, enum om
device_resident ( variable-list )
firstprivate (variable-list )
link ( variable-list )
- present ( variable-list )
- present_or_copy ( variable-list )
- pcopy ( variable-list )
- present_or_copyin ( variable-list )
- pcopyin ( variable-list )
- present_or_copyout ( variable-list )
- pcopyout ( variable-list )
- present_or_create ( variable-list )
- pcreate ( variable-list ) */
+ present ( variable-list ) */
static tree
cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
@@ -30196,16 +30190,16 @@ cp_parser_oacc_data_clause (cp_parser *parser, pra
switch (c_kind)
{
case PRAGMA_OACC_CLAUSE_COPY:
- kind = GOMP_MAP_FORCE_TOFROM;
+ kind = GOMP_MAP_TOFROM;
break;
case PRAGMA_OACC_CLAUSE_COPYIN:
- kind = GOMP_MAP_FORCE_TO;
+ kind = GOMP_MAP_TO;
break;
case PRAGMA_OACC_CLAUSE_COPYOUT:
- kind = GOMP_MAP_FORCE_FROM;
+ kind = GOMP_MAP_FROM;
break;
case PRAGMA_OACC_CLAUSE_CREATE:
- kind = GOMP_MAP_FORCE_ALLOC;
+ kind = GOMP_MAP_ALLOC;
break;
case PRAGMA_OACC_CLAUSE_DELETE:
kind = GOMP_MAP_DELETE;
@@ -30225,18 +30219,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pra
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- kind = GOMP_MAP_TOFROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- kind = GOMP_MAP_TO;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- kind = GOMP_MAP_FROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- kind = GOMP_MAP_ALLOC;
- break;
default:
gcc_unreachable ();
}
@@ -32438,22 +32420,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "present";
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copy";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyin";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyout";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_create";
- break;
case PRAGMA_OACC_CLAUSE_PRIVATE:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE,
clauses);
@@ -35387,11 +35353,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
static tree
cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
@@ -35446,11 +35408,7 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_to
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
static tree
cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
@@ -35482,8 +35440,8 @@ cp_parser_oacc_declare (cp_parser *parser, cp_toke
switch (OMP_CLAUSE_MAP_KIND (t))
{
case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
break;
@@ -35592,8 +35550,6 @@ cp_parser_oacc_declare (cp_parser *parser, cp_toke
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
#define OACC_EXIT_DATA_CLAUSE_MASK \
@@ -35724,11 +35680,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
#define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
@@ -35748,10 +35700,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
Index: gcc/testsuite/c-c++-common/goacc/default-4.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/default-4.c (revision 246808)
+++ gcc/testsuite/c-c++-common/goacc/default-4.c (working copy)
@@ -8,7 +8,7 @@ void f1 ()
float f1_b[2];
#pragma acc data copyin (f1_a) copyout (f1_b)
- /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f1_b \[^\\)\]+\\) map\\(force_to:f1_a" 1 "gimple" } } */
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f1_b \[^\\)\]+\\) map\\(to:f1_a" 1 "gimple" } } */
{
#pragma acc kernels
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */
@@ -29,7 +29,7 @@ void f2 ()
float f2_b[2];
#pragma acc data copyin (f2_a) copyout (f2_b)
- /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f2_b \[^\\)\]+\\) map\\(force_to:f2_a" 1 "gimple" } } */
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2_b \[^\\)\]+\\) map\\(to:f2_a" 1 "gimple" } } */
{
#pragma acc kernels default (none)
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */
@@ -50,7 +50,7 @@ void f3 ()
float f3_b[2];
#pragma acc data copyin (f3_a) copyout (f3_b)
- /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3_b \[^\\)\]+\\) map\\(to:f3_a" 1 "gimple" } } */
{
#pragma acc kernels default (present)
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
Index: gcc/testsuite/c-c++-common/goacc/declare-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/declare-1.c (revision 246808)
+++ gcc/testsuite/c-c++-common/goacc/declare-1.c (working copy)
@@ -19,6 +19,12 @@ int v4;
int v5, v6, v7, v8;
#pragma acc declare create(v5, v6) copyin(v7, v8)
+int v9;
+#pragma acc declare present_or_copyin(v9)
+
+int v10;
+#pragma acc declare present_or_create(v10)
+
void
f (void)
{
@@ -49,6 +55,12 @@ f (void)
extern int ve4;
#pragma acc declare link(ve4)
+ extern int ve5;
+#pragma acc declare present_or_copyin(ve5)
+
+ extern int ve6;
+#pragma acc declare present_or_create(ve6)
+
int va5;
#pragma acc declare copy(va5)
Index: gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/acc-data-chain.c (revision 246808)
+++ gcc/testsuite/c-c++-common/goacc/acc-data-chain.c (working copy)
@@ -20,5 +20,5 @@ int main(int argc, char *argv[])
return 0;
}
-// { dg-final { scan-tree-dump-times "omp target oacc_data map.force_from:b.0. .len: 400.. map.force_to:a.0. .len: 400.." 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "omp target oacc_data map.from:b.0. .len: 400.. map.to:a.0. .len: 400.." 1 "gimple" } }
// { dg-final { scan-tree-dump-times "omp target oacc_parallel map.force_present:b.0. .len: 400.. map.firstprivate:b .pointer assign, bias: 0.. map.force_present:a.0. .len: 400.. map.firstprivate:a .pointer assign, bias: 0.." 1 "gimple" } }
Index: gcc/testsuite/c-c++-common/goacc/declare-2.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/declare-2.c (revision 246808)
+++ gcc/testsuite/c-c++-common/goacc/declare-2.c (working copy)
@@ -29,14 +29,8 @@ int v6;
#pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */
int v7;
-#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */
+#pragma acc declare present_or_copyout(v7) /* { dg-error "at file scope" } */
-int v8;
-#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */
-
-int v9;
-#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */
-
int va10;
#pragma acc declare create (va10)
#pragma acc declare link (va10) /* { dg-error "more than once" } */
@@ -67,13 +61,7 @@ f (void)
#pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */
extern int ve4;
-#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */
+#pragma acc declare present_or_copyout(ve4) /* { dg-error "invalid use of" } */
- extern int ve5;
-#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */
-
- extern int ve6;
-#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */
-
-#pragma acc declare present (v9) /* { dg-error "invalid use of" } */
+#pragma acc declare present (v2) /* { dg-error "invalid use of" } */
}
Index: gcc/testsuite/gfortran.dg/goacc/declare-2.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/declare-2.f95 (revision 246808)
+++ gcc/testsuite/gfortran.dg/goacc/declare-2.f95 (working copy)
@@ -11,9 +11,9 @@ subroutine asubr (b)
!$acc declare copyout (b) ! { dg-error "Invalid clause in module" }
!$acc declare present (b) ! { dg-error "Invalid clause in module" }
!$acc declare present_or_copy (b) ! { dg-error "Invalid clause in module" }
- !$acc declare present_or_copyin (b) ! { dg-error "Invalid clause in module" }
+ !$acc declare present_or_copyin (b) ! { dg-error "present on multiple clauses" }
!$acc declare present_or_copyout (b) ! { dg-error "Invalid clause in module" }
- !$acc declare present_or_create (b) ! { dg-error "Invalid clause in module" }
+ !$acc declare present_or_create (b) ! { dg-error "present on multiple clauses" }
!$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" }
!$acc declare create (b) copyin (b) ! { dg-error "present on multiple clauses" }
Index: gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 (revision 246808)
+++ gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 (working copy)
@@ -27,10 +27,10 @@ end program test
! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } }
! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
Index: gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 (revision 246808)
+++ gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 (working copy)
@@ -166,5 +166,5 @@ end subroutine test
! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" { xfail *-*-* } } }
! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" { xfail *-*-* } } }
! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.force_tofrom:y" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.tofrom:y" 2 "gimple" } }
! { dg-final { scan-tree-dump-times "acc loop private.i. reduction..:y." 2 "gimple" { xfail *-*-* } } }
Index: gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 (revision 246808)
+++ gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 (working copy)
@@ -38,9 +38,7 @@ program test
!$acc end parallel
end program test
-! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } }
-! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } }
-! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.tofrom:v1" 9 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.tofrom:v2" 9 "gimple" } }
! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } }
Index: gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 (revision 246808)
+++ gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 (working copy)
@@ -17,10 +17,10 @@ end program test
! { dg-final { scan-tree-dump-times "if" 1 "original" } }
! { dg-final { scan-tree-dump-times "async" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
Index: gcc/testsuite/gfortran.dg/goacc/default-4.f
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/default-4.f (revision 246808)
+++ gcc/testsuite/gfortran.dg/goacc/default-4.f (working copy)
@@ -8,7 +8,7 @@
REAL, DIMENSION (2) :: F1_B
!$ACC DATA COPYIN (F1_A) COPYOUT (F1_B)
-! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f1_a \[^\\)\]+\\) map\\(force_from:f1_b" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f1_a \[^\\)\]+\\) map\\(from:f1_b" 1 "gimple" } }
!$ACC KERNELS
! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } }
F1_B(1) = F1_A;
@@ -26,7 +26,7 @@
REAL, DIMENSION (2) :: F2_B
!$ACC DATA COPYIN (F2_A) COPYOUT (F2_B)
-! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f2_a \[^\\)\]+\\) map\\(force_from:f2_b" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2_a \[^\\)\]+\\) map\\(from:f2_b" 1 "gimple" } }
!$ACC KERNELS DEFAULT (NONE)
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } }
F2_B(1) = F2_A;
@@ -44,7 +44,7 @@
REAL, DIMENSION (2) :: F3_B
!$ACC DATA COPYIN (F3_A) COPYOUT (F3_B)
-! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3_a \[^\\)\]+\\) map\\(from:f3_b" 1 "gimple" } }
!$ACC KERNELS DEFAULT (PRESENT)
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
F3_B(1) = F3_A;
Index: gcc/testsuite/gfortran.dg/goacc/data-tree.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/data-tree.f95 (revision 246808)
+++ gcc/testsuite/gfortran.dg/goacc/data-tree.f95 (working copy)
@@ -15,10 +15,10 @@ end program test
! { dg-final { scan-tree-dump-times "pragma acc data" 1 "original" } }
! { dg-final { scan-tree-dump-times "if" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c (revision 246808)
+++ libgomp/oacc-parallel.c (working copy)
@@ -492,6 +492,7 @@ GOACC_enter_exit_data (int device, size_t mapnum,
}
if (kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_FROM
|| kind == GOMP_MAP_FORCE_FROM
|| kind == GOMP_MAP_DECLARE_DEALLOCATE)
break;
@@ -573,6 +574,7 @@ GOACC_enter_exit_data (int device, size_t mapnum,
if (acc_is_present (hostaddrs[i], sizes[i]))
acc_delete (hostaddrs[i], sizes[i]);
break;
+ case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_FROM:
acc_copyout (hostaddrs[i], sizes[i]);
break;
@@ -589,9 +591,9 @@ GOACC_enter_exit_data (int device, size_t mapnum,
&sizes[i], &kinds[i]);
else if (acc_is_present (hostaddrs[i], sizes[i]))
{
- gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
- == GOMP_MAP_FORCE_FROM, async,
- pointer);
+ bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
+ || kind == GOMP_MAP_FROM);
+ gomp_acc_remove_pointer (hostaddrs[i], copyfrom, async, pointer);
/* See the above comment. */
}
i += pointer - 1;
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f (working copy)
@@ -1,16 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
- CALL ACC_PRESENT_OR_COPYIN (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_COPYIN (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f (working copy)
@@ -1,18 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
-
- INTEGER I
-
-!$ACC DATA CREATE (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC PARALLEL COPYIN (I)
- I = 0
-!$ACC END PARALLEL
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f (working copy)
@@ -1,18 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
- CALL ACC_COPYIN (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC DATA COPY (I)
- I = 0
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f (working copy)
@@ -1,16 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
-!$ACC ENTER DATA CREATE (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_COPYIN (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f (working copy)
@@ -1,18 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
-
- INTEGER I
-
-!$ACC DATA PRESENT_OR_COPY (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC DATA COPYOUT (I)
- I = 0
-!$ACC END DATA
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f (working copy)
@@ -1,16 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
- CALL ACC_PRESENT_OR_COPYIN (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC ENTER DATA CREATE (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f (working copy)
@@ -1,17 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
-!$ACC DATA PRESENT_OR_COPY (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_COPYIN (I)
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f (working copy)
@@ -1,16 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
-!$ACC ENTER DATA CREATE (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_CREATE (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c (working copy)
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
- acc_present_or_copyin (&i, sizeof i);
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyin (&i, sizeof i);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c (working copy)
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc enter data create (i)
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyin (&i, sizeof i);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c (working copy)
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
- acc_present_or_copyin (&i, sizeof i);
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc enter data create (i)
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c (working copy)
@@ -1,4 +1,4 @@
-/* Test if duplicate data mappings with acc_copy_in. */
+/* Test if acc_copyin has present_or_ behavior. */
/* { dg-do run { target openacc_nvidia_accel_selected } } */
@@ -31,5 +31,3 @@ main (int argc, char **argv)
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c (working copy)
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc enter data create (i)
- fprintf (stderr, "CheCKpOInT\n");
- acc_create (&i, sizeof i);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c (working copy)
@@ -32,5 +32,3 @@ main (int argc, char **argv)
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c (working copy)
@@ -1,22 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc data create (i)
- {
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc parallel copyin (i)
- ++i;
- }
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c (working copy)
@@ -1,22 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
- acc_copyin (&i, sizeof i);
-
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc data copy (i)
- ++i;
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c (working copy)
@@ -1,22 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc data present_or_copy (i)
- {
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc data copyout (i)
- ++i;
- }
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
-/* { dg-shouldfail "" } */
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c (revision 246808)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c (working copy)
@@ -1,22 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc data present_or_copy (i)
- {
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyin (&i, sizeof i);
- }
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
Index: libgomp/oacc-mem.c
===================================================================
--- libgomp/oacc-mem.c (revision 246808)
+++ libgomp/oacc-mem.c (working copy)
@@ -524,39 +524,49 @@ present_create_copy (unsigned f, void *h, size_t s
void *
acc_create (void *h, size_t s)
{
- return present_create_copy (FLAG_CREATE, h, s, acc_async_sync);
+ return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, acc_async_sync);
}
void
acc_create_async (void *h, size_t s, int async)
{
- present_create_copy (FLAG_CREATE, h, s, async);
+ present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, async);
}
+#ifdef HAVE_ATTRIBUTE_ALIAS
+extern void *acc_present_or_create (void *, size_t)
+ __attribute__((alias ("acc_create")));
+#else
void *
+acc_present_or_create (void *h, size_t s)
+{
+ return acc_create (h, s);
+}
+#endif
+
+void *
acc_copyin (void *h, size_t s)
{
- return present_create_copy (FLAG_CREATE | FLAG_COPY, h, s, acc_async_sync);
+ return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s,
+ acc_async_sync);
}
void
acc_copyin_async (void *h, size_t s, int async)
{
- present_create_copy (FLAG_CREATE | FLAG_COPY, h, s, async);
+ present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, async);
}
+#ifdef HAVE_ATTRIBUTE_ALIAS
+extern void *acc_present_or_copyin (void *, size_t)
+ __attribute__((alias ("acc_copyin")));
+#else
void *
-acc_present_or_create (void *h, size_t s)
-{
- return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, acc_async_sync);
-}
-
-void *
acc_present_or_copyin (void *h, size_t s)
{
- return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s,
- acc_async_sync);
+ return acc_copyin (h, s);
}
+#endif
#define FLAG_COPYOUT (1 << 0)