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 <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@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
