This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[PATCH, og8] Add OpenACC 2.6 `no_create' clause support
- From: "Maciej W. Rozycki" <macro at codesourcery dot com>
- To: <gcc-patches at gcc dot gnu dot org>
- Cc: Julian Brown <julian at codesourcery dot com>, Thomas Schwinge <thomas at codesourcery dot com>, Chung-Lin Tang <cltang at codesourcery dot com>, Jakub Jelinek <jakub at redhat dot com>, Catherine Moore <clm at codesourcery dot com>
- Date: Wed, 19 Dec 2018 21:30:19 +0000
- Subject: [PATCH, og8] Add OpenACC 2.6 `no_create' clause support
The clause makes any device code use the local memory address for each
of the variables specified unless the given variable is already present
on the current device.
2018-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
gcc/
* omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC.
* tree-pretty-print.c (dump_omp_clause): Likewise.
gcc/c-family/
* c-pragma.h (pragma_omp_clause): Add
PRAGMA_OACC_CLAUSE_NO_CREATE.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Support no_create.
(c_parser_oacc_data_clause): Likewise.
(c_parser_oacc_all_clauses): Likewise.
(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
(OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add
PRAGMA_OACC_CLAUSE_NO_CREATE.
* c-typeck.c (handle_omp_array_sections): Support
GOMP_MAP_NO_ALLOC.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Support no_create.
(cp_parser_oacc_data_clause): Likewise.
(cp_parser_oacc_all_clauses): Likewise.
(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
(OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add
PRAGMA_OACC_CLAUSE_NO_CREATE.
* semantics.c (handle_omp_array_sections): Support no_create.
gcc/fortran/
* gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC.
* openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE.
(gfc_match_omp_clauses): Support no_create.
(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES)
(OACC_SERIAL_CLAUSES, OACC_DATA_CLAUSES): Add
OMP_CLAUSE_NO_CREATE.
* trans-openmp.c (gfc_trans_omp_clauses_1): Support
OMP_MAP_NO_ALLOC.
include/
* gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC.
libgomp/
* target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC.
* testsuite/libgomp.oacc-c-c++-common/nocreate-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/nocreate-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/nocreate-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/nocreate-4.c: New test.
* testsuite/libgomp.oacc-fortran/nocreate-1.f90: New test.
* testsuite/libgomp.oacc-fortran/nocreate-2.f90: New test.
---
Hi,
This has passed regression-testing with the `x86_64-linux-gnu' target and
the `nvptx-none' offload target, across the `gcc', `g++', `gfortran' and
`libgomp' test suites. I will appreciate feedback and if none has been
given shortly, then I will commit this change to the og8 branch.
Maciej
---
gcc/c-family/c-pragma.h | 1
gcc/c/c-parser.c | 20 ++++
gcc/c/c-typeck.c | 1
gcc/cp/parser.c | 20 ++++
gcc/cp/semantics.c | 1
gcc/fortran/gfortran.h | 1
gcc/fortran/openmp.c | 15 ++-
gcc/fortran/trans-openmp.c | 3
gcc/omp-low.c | 2
gcc/tree-pretty-print.c | 3
include/gomp-constants.h | 2
libgomp/target.c | 53 +++++++++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c | 40 +++++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c | 28 ++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c | 38 +++++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c | 42 ++++++++++
libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 | 29 +++++++
libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 | 61 +++++++++++++++
18 files changed, 352 insertions(+), 8 deletions(-)
gcc-openacc-no-create.diff
Index: gcc-openacc-gcc-8-branch/gcc/c-family/c-pragma.h
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/c-family/c-pragma.h
+++ gcc-openacc-gcc-8-branch/gcc/c-family/c-pragma.h
@@ -147,6 +147,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
+ PRAGMA_OACC_CLAUSE_NO_CREATE,
PRAGMA_OACC_CLAUSE_NOHOST,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
Index: gcc-openacc-gcc-8-branch/gcc/c/c-parser.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/c/c-parser.c
+++ gcc-openacc-gcc-8-branch/gcc/c/c-parser.c
@@ -11315,7 +11315,9 @@ c_parser_omp_clause_name (c_parser *pars
result = PRAGMA_OMP_CLAUSE_MERGEABLE;
break;
case 'n':
- if (!strcmp ("nogroup", p))
+ if (!strcmp ("no_create", p))
+ result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+ else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("notinbranch", p))
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
@@ -11689,7 +11691,10 @@ c_parser_omp_var_list_parens (c_parser *
create ( variable-list )
delete ( variable-list )
detach ( variable-list )
- present ( variable-list ) */
+ present ( variable-list )
+
+ OpenACC 2.6:
+ no_create ( variable-list ) */
static tree
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
@@ -11731,6 +11736,9 @@ c_parser_oacc_data_clause (c_parser *par
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_NO_ALLOC;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -14194,6 +14202,10 @@ c_parser_oacc_all_clauses (c_parser *par
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "link";
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "no_create";
+ break;
case PRAGMA_OACC_CLAUSE_NOHOST:
clauses = c_parser_oacc_simple_clause (parser, here,
OMP_CLAUSE_NOHOST, clauses);
@@ -14619,6 +14631,7 @@ c_parser_oacc_cache (location_t loc, c_p
| (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_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
static tree
@@ -14968,6 +14981,7 @@ c_parser_oacc_loop (location_t loc, c_pa
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (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) \
@@ -14992,6 +15006,7 @@ c_parser_oacc_loop (location_t loc, c_pa
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
@@ -15019,6 +15034,7 @@ c_parser_oacc_loop (location_t loc, c_pa
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
Index: gcc-openacc-gcc-8-branch/gcc/c/c-typeck.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/c/c-typeck.c
+++ gcc-openacc-gcc-8-branch/gcc/c/c-typeck.c
@@ -12978,6 +12978,7 @@ handle_omp_array_sections (tree c, enum
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_NO_ALLOC:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
Index: gcc-openacc-gcc-8-branch/gcc/cp/parser.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/cp/parser.c
+++ gcc-openacc-gcc-8-branch/gcc/cp/parser.c
@@ -31353,7 +31353,9 @@ cp_parser_omp_clause_name (cp_parser *pa
result = PRAGMA_OMP_CLAUSE_MERGEABLE;
break;
case 'n':
- if (!strcmp ("nogroup", p))
+ if (!strcmp ("no_create", p))
+ result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+ else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nohost", p))
result = PRAGMA_OACC_CLAUSE_NOHOST;
@@ -31694,7 +31696,10 @@ cp_parser_omp_var_list (cp_parser *parse
create ( variable-list )
delete ( variable-list )
detach ( variable-list )
- present ( variable-list ) */
+ present ( variable-list )
+
+ OpenACC 2.6:
+ no_create ( variable-list ) */
static tree
cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
@@ -31736,6 +31741,9 @@ cp_parser_oacc_data_clause (cp_parser *p
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_NO_ALLOC;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -33964,6 +33972,10 @@ cp_parser_oacc_all_clauses (cp_parser *p
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "link";
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "no_create";
+ break;
case PRAGMA_OACC_CLAUSE_NOHOST:
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_NOHOST,
clauses, here);
@@ -36936,6 +36948,7 @@ cp_parser_oacc_cache (cp_parser *parser,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
static tree
@@ -37272,6 +37285,7 @@ cp_parser_oacc_loop (cp_parser *parser,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (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) \
@@ -37297,6 +37311,7 @@ cp_parser_oacc_loop (cp_parser *parser,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (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) \
@@ -37323,6 +37338,7 @@ cp_parser_oacc_loop (cp_parser *parser,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
Index: gcc-openacc-gcc-8-branch/gcc/cp/semantics.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/cp/semantics.c
+++ gcc-openacc-gcc-8-branch/gcc/cp/semantics.c
@@ -5096,6 +5096,7 @@ handle_omp_array_sections (tree c, enum
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_NO_ALLOC:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
Index: gcc-openacc-gcc-8-branch/gcc/fortran/gfortran.h
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/fortran/gfortran.h
+++ gcc-openacc-gcc-8-branch/gcc/fortran/gfortran.h
@@ -1184,6 +1184,7 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_NO_ALLOC,
OMP_MAP_ATTACH,
OMP_MAP_TO,
OMP_MAP_FROM,
Index: gcc-openacc-gcc-8-branch/gcc/fortran/openmp.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/fortran/openmp.c
+++ gcc-openacc-gcc-8-branch/gcc/fortran/openmp.c
@@ -818,6 +818,7 @@ enum omp_mask2
OMP_CLAUSE_COPY,
OMP_CLAUSE_COPYOUT,
OMP_CLAUSE_CREATE,
+ OMP_CLAUSE_NO_CREATE,
OMP_CLAUSE_PRESENT,
OMP_CLAUSE_DEVICEPTR,
OMP_CLAUSE_GANG,
@@ -1559,6 +1560,12 @@ gfc_match_omp_clauses (gfc_omp_clauses *
}
break;
case 'n':
+ if ((mask & OMP_CLAUSE_NO_CREATE)
+ && gfc_match ("no_create ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_NO_ALLOC, true,
+ allow_derived))
+ continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
&& gfc_match ("nogroup") == MATCH_YES)
@@ -2070,7 +2077,7 @@ gfc_match_omp_clauses (gfc_omp_clauses *
| OMP_CLAUSE_IF \
| OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH)
@@ -2081,7 +2088,7 @@ gfc_match_omp_clauses (gfc_omp_clauses *
| OMP_CLAUSE_DEVICE_TYPE \
| OMP_CLAUSE_IF \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH)
#define OACC_SERIAL_CLAUSES \
@@ -2090,14 +2097,14 @@ gfc_match_omp_clauses (gfc_omp_clauses *
| OMP_CLAUSE_IF \
| OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_ATTACH)
#define OACC_HOST_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_USE_DEVICE))
Index: gcc-openacc-gcc-8-branch/gcc/fortran/trans-openmp.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/fortran/trans-openmp.c
+++ gcc-openacc-gcc-8-branch/gcc/fortran/trans-openmp.c
@@ -2348,6 +2348,9 @@ gfc_trans_omp_clauses_1 (stmtblock_t *bl
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
+ case OMP_MAP_NO_ALLOC:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_NO_ALLOC);
+ break;
case OMP_MAP_ATTACH:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
break;
Index: gcc-openacc-gcc-8-branch/gcc/omp-low.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/omp-low.c
+++ gcc-openacc-gcc-8-branch/gcc/omp-low.c
@@ -8184,6 +8184,7 @@ lower_omp_target (gimple_stmt_iterator *
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
break;
+ case GOMP_MAP_NO_ALLOC:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
@@ -8681,6 +8682,7 @@ lower_omp_target (gimple_stmt_iterator *
switch (tkind)
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_NO_ALLOC:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
Index: gcc-openacc-gcc-8-branch/gcc/tree-pretty-print.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/gcc/tree-pretty-print.c
+++ gcc-openacc-gcc-8-branch/gcc/tree-pretty-print.c
@@ -684,6 +684,9 @@ dump_omp_clause (pretty_printer *pp, tre
case GOMP_MAP_POINTER:
pp_string (pp, "alloc");
break;
+ case GOMP_MAP_NO_ALLOC:
+ pp_string (pp, "no_alloc");
+ break;
case GOMP_MAP_TO:
case GOMP_MAP_TO_PSET:
pp_string (pp, "to");
Index: gcc-openacc-gcc-8-branch/include/gomp-constants.h
===================================================================
--- gcc-openacc-gcc-8-branch.orig/include/gomp-constants.h
+++ gcc-openacc-gcc-8-branch/include/gomp-constants.h
@@ -80,6 +80,8 @@ enum gomp_map_kind
GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1),
/* OpenACC link. */
GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2),
+ /* Use device data if present, fall back to host address otherwise. */
+ GOMP_MAP_NO_ALLOC = (GOMP_MAP_FLAG_SPECIAL_1 | 3),
/* Allocate. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
Index: gcc-openacc-gcc-8-branch/libgomp/target.c
===================================================================
--- gcc-openacc-gcc-8-branch.orig/libgomp/target.c
+++ gcc-openacc-gcc-8-branch/libgomp/target.c
@@ -1212,6 +1212,12 @@ gomp_map_vars_async (struct gomp_device_
has_firstprivate = true;
continue;
}
+ else if ((kind & typemask) == GOMP_MAP_NO_ALLOC)
+ {
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = 0;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask)
&& (kind & typemask) != GOMP_MAP_ATTACH)
@@ -1496,6 +1502,53 @@ gomp_map_vars_async (struct gomp_device_
cbufp);
continue;
}
+ case GOMP_MAP_NO_ALLOC:
+ {
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizes[i];
+ splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ if (n != NULL)
+ {
+ tgt->list[i].key = n;
+ tgt->list[i].offset = cur_node.host_start - n->host_start;
+ tgt->list[i].length = n->host_end - n->host_start;
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].do_detach = false;
+ n->refcount++;
+ }
+ else
+ {
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = OFFSET_INLINED;
+ tgt->list[i].length = sizes[i];
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].do_detach = false;
+ if (i + 1 < mapnum)
+ {
+ int kind2 = get_kind (short_mapkind, kinds, i + 1);
+ switch (kind2 & typemask)
+ {
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_POINTER:
+ /* The data is not present but we have an attach
+ or pointer clause next. Skip over it. */
+ i++;
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = OFFSET_INLINED;
+ tgt->list[i].length = sizes[i];
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].do_detach = false;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ continue;
+ }
default:
break;
}
Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c
===================================================================
--- /dev/null
+++ gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c
@@ -0,0 +1,40 @@
+/* Test no_create clause when data is present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr;
+
+ acc_copyin (arr, N * sizeof (*arr));
+
+ #pragma acc parallel no_create(arr[0:N]) copyout(devptr)
+ {
+ devptr = &arr[2];
+ }
+
+#if !ACC_MEM_SHARED
+ if (acc_hostptr (devptr) != (void *) &arr[2])
+ __builtin_abort ();
+#endif
+
+ acc_delete (arr, N * sizeof (*arr));
+
+#if ACC_MEM_SHARED
+ if (&arr[2] != devptr)
+ __builtin_abort ();
+#else
+ if (&arr[2] == devptr)
+ __builtin_abort ();
+#endif
+
+ free (arr);
+
+ return 0;
+}
Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c
===================================================================
--- /dev/null
+++ gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c
@@ -0,0 +1,28 @@
+/* Test no_create clause when data is not present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr;
+
+ #pragma acc data no_create(arr[0:N])
+ {
+ #pragma acc parallel copyout(devptr)
+ {
+ devptr = &arr[2];
+ }
+ }
+
+ if (devptr != &arr[2])
+ __builtin_abort ();
+
+ free (arr);
+
+ return 0;
+}
Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c
===================================================================
--- /dev/null
+++ gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c
@@ -0,0 +1,38 @@
+/* Test no_create clause with attach/detach when data is not present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+typedef struct {
+ int x;
+ int *y;
+} mystruct;
+
+int
+main (int argc, char *argv[])
+{
+ int *devptr;
+ mystruct s;
+
+ s.x = 5;
+ s.y = (int *) malloc (N * sizeof (int));
+
+ #pragma acc data copyin(s)
+ {
+ #pragma acc parallel no_create(s.y[0:N]) copyout(devptr)
+ {
+ devptr = &s.y[2];
+ }
+ }
+
+ if (devptr != &s.y[2])
+ __builtin_abort ();
+
+ free (s.y);
+
+ return 0;
+}
Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c
===================================================================
--- /dev/null
+++ gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c
@@ -0,0 +1,42 @@
+/* Test no_create clause with attach/detach when data is present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+typedef struct {
+ int x;
+ int *y;
+} mystruct;
+
+int
+main (int argc, char *argv[])
+{
+ int *devptr;
+ mystruct s;
+
+ s.x = 5;
+ s.y = (int *) malloc (N * sizeof (int));
+
+ #pragma acc data copyin(s)
+ {
+ #pragma acc enter data copyin(s.y[0:N])
+
+ #pragma acc parallel no_create(s.y[0:N]) copyout(devptr)
+ {
+ devptr = &s.y[2];
+ }
+ }
+
+ if (devptr != acc_deviceptr (&s.y[2]))
+ __builtin_abort ();
+
+ #pragma acc exit data delete(s.y[0:N])
+
+ free (s.y);
+
+ return 0;
+}
Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90
===================================================================
--- /dev/null
+++ gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90
@@ -0,0 +1,29 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Test no_create clause with data construct when data is present/not present.
+
+program nocreate
+ use openacc
+ implicit none
+ integer, parameter :: n = 512
+ integer :: myarr(n)
+ integer i
+
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc data no_create (myarr)
+ if (acc_is_present (myarr)) stop 1
+ !$acc end data
+
+ !$acc enter data copyin (myarr)
+ !$acc data no_create (myarr)
+ if (acc_is_present (myarr) .eqv. .false.) stop 2
+ !$acc end data
+ !$acc exit data copyout (myarr)
+
+ do i = 1, n
+ if (myarr(i) .ne. 0) stop 3
+ end do
+end program nocreate
Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90
===================================================================
--- /dev/null
+++ gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90
@@ -0,0 +1,61 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Test no_create clause with data/parallel constructs.
+
+program nocreate
+ use openacc
+ implicit none
+ integer, parameter :: n = 512
+ integer :: myarr(n)
+ integer i
+
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ call do_on_target(myarr, n)
+
+ do i = 1, n
+ if (myarr(i) .ne. i) stop 1
+ end do
+
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc enter data copyin(myarr)
+ call do_on_target(myarr, n)
+ !$acc exit data copyout(myarr)
+
+ do i = 1, n
+ if (myarr(i) .ne. i * 2) stop 2
+ end do
+end program nocreate
+
+subroutine do_on_target (arr, n)
+ use openacc
+ implicit none
+ integer :: n, arr(n)
+ integer :: i
+
+!$acc data no_create (arr)
+
+if (acc_is_present(arr)) then
+ ! The no_create clause is meant for partially shared-memory machines. This
+ ! test is written to work on non-shared-memory machines, though this is not
+ ! necessarily a useful way to use the no_create clause in practice.
+
+ !$acc parallel loop no_create (arr)
+ do i = 1, n
+ arr(i) = i * 2
+ end do
+ !$acc end parallel loop
+else
+ do i = 1, n
+ arr(i) = i
+ end do
+end if
+
+!$acc end data
+
+end subroutine do_on_target