* [PATCH, og8] Add OpenACC 2.6 `no_create' clause support
@ 2018-12-19 21:31 ` Maciej W. Rozycki
2018-12-20 14:36 ` Maciej W. Rozycki
2019-10-24 13:26 ` [Patch] Add OpenACC 2.6's no_create Tobias Burnus
0 siblings, 2 replies; 12+ messages in thread
From: Maciej W. Rozycki @ 2018-12-19 21:31 UTC (permalink / raw)
To: gcc-patches
Cc: Julian Brown, Thomas Schwinge, Chung-Lin Tang, Jakub Jelinek,
Catherine Moore
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
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [PATCH, og8] Add OpenACC 2.6 `no_create' clause support
2018-12-19 21:31 ` [PATCH, og8] Add OpenACC 2.6 `no_create' clause support Maciej W. Rozycki
@ 2018-12-20 14:36 ` Maciej W. Rozycki
2019-10-24 13:26 ` [Patch] Add OpenACC 2.6's no_create Tobias Burnus
1 sibling, 0 replies; 12+ messages in thread
From: Maciej W. Rozycki @ 2018-12-20 14:36 UTC (permalink / raw)
To: gcc-patches
Cc: Julian Brown, Thomas Schwinge, Chung-Lin Tang, Jakub Jelinek,
Catherine Moore
On Wed, 19 Dec 2018, Maciej W. Rozycki wrote:
> 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.
I have committed this change now.
Maciej
^ permalink raw reply [flat|nested] 12+ messages in thread
* [Patch] Add OpenACC 2.6's no_create
@ 2019-10-24 13:26 ` Tobias Burnus
2019-11-05 23:49 ` Thomas Schwinge
0 siblings, 1 reply; 12+ messages in thread
From: Tobias Burnus @ 2019-10-24 13:26 UTC (permalink / raw)
To: gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 1709 bytes --]
The clause (new in OpenACC 2.6) 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. – Or in words of
OpenACC 2.7 (in Sect. 2.7.9 no_create clause):
"The no_create clause may appear on structured data and compute
constructs." / "For each var in varlist, if var is in shared memory, no
action is taken; if var is not in shared memory, the no_create clause
behaves as follows:" [digest: if present, update present count, if
pointer attach/detach; if not not present, device-local memory used.]
"The restrictions regarding subarrays in the present clause apply to
this clause."
Note: The "no_create" maps to the (new) GOMP_MAP_NO_ALLOC in the middle
end – and all action in libgomp/target.c but only applies to
GOMP_MAP_NO_ALLOC; hence, the code should only affect OpenACC.
OK for the trunk?
Cheers,
Tobias
PS: This patch is a re-diffed version of the OG9/OG8 version; as some
other features are not yet on trunk, it misses a test case for
"no_create(s.y…)" (i.e. the struct component-ref;
libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-{3,4}.c); trunk
also lacks 'acc serial' and, hence, the attach patch lacks the
OACC_SERIAL_CLAUSE_MASK updates – and gfc_match_omp_map_clause needs
later to be updated for the allow_derived and allow_common arguments.
Furthermore, some 'do_detach = false' are missing in libgomp/target.c as
they do not yet exist on trunk, either.
The openacc-gcc-9 /…-8 branch patch is commit
8e74c2ec2b90819c995444370e742864a685209f of Dec 20, 2018. It has been
posted as https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01418.html
[-- Attachment #2: openacc_no_create2.diff --]
[-- Type: text/x-patch, Size: 20946 bytes --]
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.
2019-10-24 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@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): 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_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-fortran/nocreate-1.f90: New test.
* testsuite/libgomp.oacc-fortran/nocreate-2.f90: New test.
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index e0aa774555a..da6cfdb8b98 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -153,6 +153,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_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 7618a46c8bc..1004a2e5579 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11833,7 +11833,9 @@ c_parser_omp_clause_name (c_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -12296,7 +12298,10 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -12332,6 +12337,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
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;
@@ -15130,6 +15138,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
clauses = c_parser_oacc_single_int_clause (parser,
OMP_CLAUSE_NUM_GANGS,
@@ -15598,6 +15610,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
| (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
@@ -15925,6 +15938,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -15940,6 +15954,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index c7339509bd1..3be6b654598 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13409,6 +13409,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
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:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 3857fe47d67..8d7de8bc33b 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -33019,7 +33019,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -33385,7 +33387,10 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -33421,6 +33426,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
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;
@@ -35983,6 +35991,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@@ -39788,6 +39800,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
| (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
@@ -40105,6 +40118,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -40119,8 +40133,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
+ | (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) \
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 59def3170ab..db7cac82312 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5291,6 +5291,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
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:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 920acdafc6b..5c930097a66 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1191,6 +1191,7 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_NO_ALLOC,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 5c91fcdfd31..00575fd37aa 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -807,6 +807,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,
@@ -1444,6 +1445,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
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))
+ continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
&& gfc_match ("nogroup") == MATCH_YES)
@@ -1954,19 +1960,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
(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_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#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_NO_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 \
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index dad11a24430..979f83c234e 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2351,6 +2351,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
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_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 279b6ef893a..8cf43852b2a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11315,6 +11315,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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:
@@ -11724,6 +11725,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
switch (tkind)
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_NO_ALLOC:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 53b3f55a3e6..51cfa837c91 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -788,6 +788,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
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");
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 82e9094c934..e9f5441d9da 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -75,6 +75,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),
/* Do not map, copy bits for firstprivate instead. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
diff --git a/libgomp/target.c b/libgomp/target.c
index 84d6daa76ca..cc79bb4dd06 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -667,6 +667,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
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))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -892,6 +898,49 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
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;
+ 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;
+ if (i + 1 < mapnum)
+ {
+ int kind2 = get_kind (short_mapkind, kinds, i + 1);
+ switch (kind2 & typemask)
+ {
+ 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;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ continue;
+ }
default:
break;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c
new file mode 100644
index 00000000000..c7a1bd9c015
--- /dev/null
+++ b/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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c
new file mode 100644
index 00000000000..2964a40b217
--- /dev/null
+++ b/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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90
new file mode 100644
index 00000000000..f048355d7df
--- /dev/null
+++ b/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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90
new file mode 100644
index 00000000000..34444ecf5b0
--- /dev/null
+++ b/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
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-10-24 13:26 ` [Patch] Add OpenACC 2.6's no_create Tobias Burnus
@ 2019-11-05 23:49 ` Thomas Schwinge
2019-11-06 12:43 ` Thomas Schwinge
0 siblings, 1 reply; 12+ messages in thread
From: Thomas Schwinge @ 2019-11-05 23:49 UTC (permalink / raw)
To: Tobias Burnus, Jakub Jelinek; +Cc: gcc-patches, fortran, Julian Brown
[-- Attachment #1.1: Type: text/plain, Size: 6084 bytes --]
Hi Tobias!
On 2019-10-24T14:47:58+0200, Tobias Burnus <Tobias_Burnus@mentor.com> wrote:
> The clause (new in OpenACC 2.6) 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. – Or in words of
> OpenACC 2.7 (in Sect. 2.7.9 no_create clause):
>
> "The no_create clause may appear on structured data and compute
> constructs." / "For each var in varlist, if var is in shared memory, no
> action is taken; if var is not in shared memory, the no_create clause
> behaves as follows:" [digest: if present, update present count, if
> pointer attach/detach; if not not present, device-local memory used.]
s%not not%not%
s%device-local%local%
> "The restrictions regarding subarrays in the present clause apply to
> this clause."
> Note: The "no_create" maps to the (new) GOMP_MAP_NO_ALLOC in the middle
> end – and all action in libgomp/target.c but only applies to
> GOMP_MAP_NO_ALLOC; hence, the code should only affect OpenACC.
Not sure if 'GOMP_MAP_NO_ALLOC' is the most descriptive name. ;-)
I understand 'no_create' to mean 'present' in combination with an
'if_present' flag that is available as a clause for some (other) OpenACC
directives, correct?
So, how about naming this 'GOMP_MAP_IF_PRESENT' instead of
'GOMP_MAP_NO_ALLOC'? (Jakub?) (But I don't care too much, so if there's
a good reason to prefer 'GOMP_MAP_NO_ALLOC', then that's fine, too.)
Ah, I just found that Julian (CCed for your information) internally had
proposed 'GOMP_MAP_MAYBE_PRESENT', which seems like another good option
indeed.
For reference:
> --- a/include/gomp-constants.h
> +++ b/include/gomp-constants.h
> @@ -75,6 +75,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),
> /* Do not map, copy bits for firstprivate instead. */
> GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
> /* Similarly, but store the value in the pointer rather than
> OK for the trunk?
To synchronize our efforts, I'm attaching an incremental WIP patch. Will
you please have a look at that, merging it in, while I continue to
review?
> PS: This patch is a re-diffed version of the OG9/OG8 version; as some
> other features are not yet on trunk, it misses a test case for
> "no_create(s.y…)" (i.e. the struct component-ref;
> libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-{3,4}.c); trunk
> also lacks 'acc serial' and, hence, the attach patch lacks the
> OACC_SERIAL_CLAUSE_MASK updates – and gfc_match_omp_map_clause needs
> later to be updated for the allow_derived and allow_common arguments.
> Furthermore, some 'do_detach = false' are missing in libgomp/target.c as
> they do not yet exist on trunk, either.
>
> The openacc-gcc-9 /…-8 branch patch is commit
> 8e74c2ec2b90819c995444370e742864a685209f of Dec 20, 2018. It has been
> posted as https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01418.html
Thanks for providing these references, that's useful.
> libgomp/
> * 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-fortran/nocreate-1.f90: New test.
> * testsuite/libgomp.oacc-fortran/nocreate-2.f90: New test.
Please rename these files to 'no_create*', as that's what the clause is
called.
..., and then:
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -667,6 +667,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> 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))
> cur_node.host_end = cur_node.host_start + sizes[i];
> @@ -892,6 +898,49 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
> + cur_node.host_start - n->host_start;
> 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;
> + 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;
> + if (i + 1 < mapnum)
> + {
> + int kind2 = get_kind (short_mapkind, kinds, i + 1);
> + switch (kind2 & typemask)
> + {
> + 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;
> + break;
> + default:
> + break;
> + }
> + }
> + }
> + continue;
> + }
> default:
> break;
> }
This I don't grok yet; see the "TODO" comments in the attached
incremental WIP patch.
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-WIP-into-Add-OpenACC-2.6-no_create-clause-support.patch --]
[-- Type: text/x-diff, Size: 12920 bytes --]
From 22ceeb89f787a6279a37d35965f82a4f5b3e3b72 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 6 Nov 2019 00:42:06 +0100
Subject: [PATCH] [WIP] into Add OpenACC 2.6 `no_create' clause support
---
gcc/fortran/openmp.c | 4 ++--
.../gfortran.dg/goacc/common-block-1.f90 | 3 +++
.../gfortran.dg/goacc/common-block-2.f90 | 3 +++
.../gfortran.dg/goacc/data-clauses.f95 | 23 ++++++++++++++++++-
gcc/testsuite/gfortran.dg/goacc/data-tree.f95 | 3 ++-
.../gfortran.dg/goacc/kernels-tree.f95 | 3 ++-
.../gfortran.dg/goacc/parallel-tree.f95 | 3 ++-
libgomp/target.c | 8 +++++++
.../libgomp.oacc-fortran/common-block-2.f90 | 4 +++-
.../libgomp.oacc-fortran/nocreate-1.f90 | 10 ++++++--
.../libgomp.oacc-fortran/nocreate-2.f90 | 16 ++++++++++---
11 files changed, 68 insertions(+), 12 deletions(-)
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 47c5cf5d422..822af5dbe7c 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1449,7 +1449,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
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))
+ OMP_MAP_NO_ALLOC, true))
continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
@@ -1969,7 +1969,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
index ea437526b46..5c162a5b884 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
@@ -51,6 +51,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v)
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v)
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
index 1ba945019f9..33c0d3f5fb4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
@@ -39,6 +39,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
index b94214e8b63..c1b3e1dec38 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
@@ -111,6 +111,27 @@ contains
!$acc end data
+ !$acc parallel no_create (tip) ! { dg-error "POINTER" }
+ !$acc end parallel
+ !$acc parallel no_create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc end parallel
+ !$acc parallel deviceptr (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copy (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyin (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyout (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end parallel
+ !$acc kernels no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end kernels
+ !$acc data no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end data
+
+
!$acc parallel present (tip) ! { dg-error "POINTER" }
!$acc end parallel
!$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }
@@ -256,4 +277,4 @@ contains
!$acc end data
end subroutine foo
-end module test
\ No newline at end of file
+end module test
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
index f16d62cce69..454417d6a05 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
@@ -7,6 +7,7 @@ program test
logical :: l = .true.
!$acc data if(l) copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end data
@@ -19,7 +20,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index a70f1e737bd..5583ffb4d04 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -8,6 +8,7 @@ program test
!$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end kernels
@@ -25,7 +26,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 2697bb79e7f..e33653bdd78 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -9,6 +9,7 @@ program test
!$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
!$acc end parallel
@@ -28,7 +29,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/libgomp/target.c b/libgomp/target.c
index 632e7020538..0338648946d 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -669,6 +669,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
else if ((kind & typemask) == GOMP_MAP_NO_ALLOC)
{
+ //TODO TS is confused. Handling this here, will inhibit 'gomp_map_vars_existing' being used a bit further below.
tgt->list[i].key = NULL;
tgt->list[i].offset = 0;
continue;
@@ -905,6 +906,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
if (n != NULL)
{
+ //TODO TS is confused. Due to the way the handling of 'GOMP_MAP_NO_ALLOC' is done in the first loop, we're here re-doing 'gomp_map_vars_existing'?
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;
@@ -914,6 +916,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
else
{
+ //TODO This is basically 'GOMP_MAP_FIRSTPRIVATE_INT' handling?
tgt->list[i].key = NULL;
tgt->list[i].offset = OFFSET_INLINED;
tgt->list[i].length = sizes[i];
@@ -925,6 +928,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
switch (kind2 & typemask)
{
case GOMP_MAP_POINTER:
+ //TODO abort();
+ //TODO This code path is exercised by 'libgomp.oacc-fortran/nocreate-2.f90'.
+ //TODO TS does not yet understand why this is needed.
+ //TODO Is this somehow similar to 'GOMP_MAP_TO_PSET' handling?
+
/* The data is not present but we have an attach
or pointer clause next. Skip over it. */
i++;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
index 018b37d00bb..ad04ca997c2 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
@@ -76,7 +76,9 @@ program main
!$acc enter data create(b)
- !$acc parallel loop pcopy(b)
+ !$acc parallel loop &
+ !$acc no_create(b) ! ... here means 'present(b)'.
+ !TODO But we get: "libgomp: cuStreamSynchronize error: an illegal memory access was encountered".
do i = 1, n
b(i) = i
end do
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90
index f048355d7df..ca9611b777c 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90
@@ -1,20 +1,26 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+! { dg-do run }
! Test no_create clause with data construct when data is present/not present.
program nocreate
use openacc
implicit none
+ logical :: shared_memory
integer, parameter :: n = 512
integer :: myarr(n)
integer i
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
do i = 1, n
myarr(i) = 0
end do
!$acc data no_create (myarr)
- if (acc_is_present (myarr)) stop 1
+ if (acc_is_present (myarr) .neqv. shared_memory) stop 1
!$acc end data
!$acc enter data copyin (myarr)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90
index 34444ecf5b0..16227b8ae22 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90
@@ -1,14 +1,20 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+! { dg-do run }
! Test no_create clause with data/parallel constructs.
program nocreate
use openacc
implicit none
+ logical :: shared_memory
integer, parameter :: n = 512
integer :: myarr(n)
integer i
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
do i = 1, n
myarr(i) = 0
end do
@@ -16,7 +22,11 @@ program nocreate
call do_on_target(myarr, n)
do i = 1, n
- if (myarr(i) .ne. i) stop 1
+ if (shared_memory) then
+ if (myarr(i) .ne. i * 2) stop 1
+ else
+ if (myarr(i) .ne. i) stop 2
+ end if
end do
do i = 1, n
@@ -28,7 +38,7 @@ program nocreate
!$acc exit data copyout(myarr)
do i = 1, n
- if (myarr(i) .ne. i * 2) stop 2
+ if (myarr(i) .ne. i * 2) stop 3
end do
end program nocreate
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-11-05 23:49 ` Thomas Schwinge
@ 2019-11-06 12:43 ` Thomas Schwinge
2019-11-15 19:12 ` Tobias Burnus
0 siblings, 1 reply; 12+ messages in thread
From: Thomas Schwinge @ 2019-11-06 12:43 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches, fortran, Julian Brown, Jakub Jelinek
[-- Attachment #1.1: Type: text/plain, Size: 990 bytes --]
Hi Tobias!
On 2019-11-06T00:47:05+0100, I wrote:
> --- a/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
> @@ -76,7 +76,9 @@ program main
>
> !$acc enter data create(b)
>
> - !$acc parallel loop pcopy(b)
> + !$acc parallel loop &
> + !$acc no_create(b) ! ... here means 'present(b)'.
> + !TODO But we get: "libgomp: cuStreamSynchronize error: an illegal memory access was encountered".
> do i = 1, n
> b(i) = i
> end do
Either I'm completely confused -- always possible ;-) -- or there's
something wrong; see the two attached test cases, not actually related to
Fortran common blocks at all. If such a basic usage of the 'no_create'
clause doesn't work...? So, again..., seems that my suspicion was right
that this patch doesn't have sufficient test coverage at all. Or, I'm
completely confused -- we still have that option, too. ;-\
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-libgomp.oacc-c-c-common-common-block-2_.c-libgomp.oa.patch --]
[-- Type: text/x-diff, Size: 2291 bytes --]
From 38fcb35dcb98b0fd709db72896455895243d8e54 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 6 Nov 2019 13:39:12 +0100
Subject: [PATCH] 'libgomp.oacc-c-c++-common/common-block-2_.c',
'libgomp.oacc-fortran/common-block-2_.f90'
---
.../common-block-2_.c | 19 +++++++++++++++
.../libgomp.oacc-fortran/common-block-2_.f90 | 23 +++++++++++++++++++
2 files changed, 42 insertions(+)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/common-block-2_.c
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/common-block-2_.f90
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/common-block-2_.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/common-block-2_.c
new file mode 100644
index 00000000000..5cf547049ab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/common-block-2_.c
@@ -0,0 +1,19 @@
+// Adapted/reduced from 'libgomp.oacc-fortran/common-block-2.f90'.
+
+int main()
+{
+#define N 100
+ float b[N];
+
+#pragma acc enter data create(b)
+
+#pragma acc parallel loop \
+ /*present(b)*/ /* ... works. */ \
+ no_create(b) /* ... here also means 'present(b)', but we get: "libgomp: cuStreamSynchronize error: an illegal memory access was encountered". */
+ for (int i = 0; i < N; ++i)
+ b[i] = i;
+
+#pragma acc exit data delete(b)
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/common-block-2_.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2_.f90
new file mode 100644
index 00000000000..f3f25869bea
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2_.f90
@@ -0,0 +1,23 @@
+! { dg-do run }
+
+! Adapted/reduced from 'libgomp.oacc-fortran/common-block-2.f90'.
+
+program main
+ implicit none
+ integer i
+ integer, parameter :: n = 100
+ real*4 b(n)
+ !common /BLOCK/ b
+
+ !$acc enter data create(b)
+
+ !$acc parallel loop &
+ !!$acc present(b) ! ... works.
+ !$acc no_create(b) ! ... here also means 'present(b)', but we get: "libgomp: cuStreamSynchronize error: an illegal memory access was encountered".
+ do i = 1, n
+ b(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc exit data delete(b)
+end program main
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-11-06 12:43 ` Thomas Schwinge
@ 2019-11-15 19:12 ` Tobias Burnus
2019-12-03 15:16 ` Thomas Schwinge
0 siblings, 1 reply; 12+ messages in thread
From: Tobias Burnus @ 2019-11-15 19:12 UTC (permalink / raw)
To: Thomas Schwinge, Tobias Burnus
Cc: gcc-patches, fortran, Julian Brown, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 2184 bytes --]
Hello Thomas, hi all,
updated version. Changes:
* Incorporate Thomas's changes
* Add no_create clause to newly added 'acc serial'
* Renamed (G)OMP_MAP_NO_ALLOC to (G)OMP_MAP_IF_PRESENT as proposed
* Make no_create.c effective by adding 'has_firstprivate = true;' to
target.c.*
(* If one tries to access c or e in the no_create-3.{c,f90} run-time
test case, plugin-nvidia rightly complains (illegal memory access),
using the created 'b' or 'd' works as tested by the test case. This
feature seems to be also broken on the OG9 branch.)
Bootstrapped and regtested without offloading and with nvptx offloading.
OK?
Tobias
PS: Remaining bits of the OG9 patch, which are not included are the
following. I think those are all attach/detach features: a test case for
"no_create(s.y…)" (i.e. the struct component-ref;
libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-{3,4}.c) and some
'do_detach = false' in libgomp/target.c. Cf. openacc-gcc-9 /…-8 branch
patch is commit 8e74c2ec2b90819c995444370e742864a685209f of Dec 20,
2018. It has been posted as
https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01418.html
On 11/6/19 1:42 PM, Thomas Schwinge wrote:
> Hi Tobias!
>
> On 2019-11-06T00:47:05+0100, I wrote:
>> --- a/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
>> +++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
>> @@ -76,7 +76,9 @@ program main
>>
>> !$acc enter data create(b)
>>
>> - !$acc parallel loop pcopy(b)
>> + !$acc parallel loop &
>> + !$acc no_create(b) ! ... here means 'present(b)'.
>> + !TODO But we get: "libgomp: cuStreamSynchronize error: an illegal memory access was encountered".
>> do i = 1, n
>> b(i) = i
>> end do
> Either I'm completely confused -- always possible ;-) -- or there's
> something wrong; see the two attached test cases, not actually related to
> Fortran common blocks at all. If such a basic usage of the 'no_create'
> clause doesn't work...? So, again..., seems that my suspicion was right
> that this patch doesn't have sufficient test coverage at all. Or, I'm
> completely confused -- we still have that option, too. ;-\
>
>
> Grüße
> Thomas
>
>
[-- Attachment #2: openacc_no_create5.diff --]
[-- Type: text/x-patch, Size: 33362 bytes --]
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.
2019-11-15 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@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): 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_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
* trans-openmp.c (gfc_trans_omp_clauses_1): Support
OMP_MAP_NO_ALLOC.
gcc/testsuite/
* gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
* gfortran.dg/goacc/common-block-1.f90: Likewise.
* gfortran.dg/goacc/data-clauses.f95: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
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-fortran/nocreate-1.f90: New test.
* testsuite/libgomp.oacc-fortran/nocreate-2.f90: New test.
* testsuite/libgomp.oacc-fortran/nocreate-3.f90: New test.
gcc/c-family/c-pragma.h | 1 +
gcc/c/c-parser.c | 20 +++++-
gcc/c/c-typeck.c | 1 +
gcc/cp/parser.c | 22 ++++++-
gcc/cp/semantics.c | 1 +
gcc/fortran/gfortran.h | 1 +
gcc/fortran/openmp.c | 28 +++++----
gcc/fortran/trans-openmp.c | 3 +
gcc/omp-low.c | 2 +
gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 | 3 +
gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 | 3 +
gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 | 21 +++++++
gcc/testsuite/gfortran.dg/goacc/data-tree.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 | 3 +-
gcc/tree-pretty-print.c | 3 +
include/gomp-constants.h | 2 +
libgomp/target.c | 50 +++++++++++++++
.../libgomp.oacc-c-c++-common/no_create-1.c | 40 ++++++++++++
.../libgomp.oacc-c-c++-common/no_create-2.c | 28 +++++++++
.../libgomp.oacc-c-c++-common/no_create-3.c | 24 ++++++++
.../testsuite/libgomp.oacc-fortran/no_create-1.f90 | 35 +++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-2.f90 | 71 ++++++++++++++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-3.f90 | 21 +++++++
24 files changed, 370 insertions(+), 19 deletions(-)
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index bfe681bb430..3754c5fda45 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -154,6 +154,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_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 5b290bf7567..d93c7b6316f 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12457,7 +12457,9 @@ c_parser_omp_clause_name (c_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -12920,7 +12922,10 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -12956,6 +12961,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -15754,6 +15762,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
clauses = c_parser_oacc_single_int_clause (parser,
OMP_CLAUSE_NUM_GANGS,
@@ -16222,6 +16234,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
| (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
@@ -16554,6 +16567,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -16569,6 +16583,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -16587,6 +16602,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 5f74a3b28d9..0f04e153ea7 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13422,6 +13422,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index c473e7fd92f..07685c8c389 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -33100,7 +33100,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -33466,7 +33468,10 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -33502,6 +33507,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -36064,6 +36072,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@@ -39869,6 +39881,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
| (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
@@ -40190,6 +40203,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -40204,8 +40218,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
+ | (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) \
@@ -40223,6 +40238,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 0ce73af5bc6..5a23b37a201 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5300,6 +5300,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index e962db59bc5..3b473854f06 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1192,6 +1192,7 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_IF_PRESENT,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index dc0521b40f0..576003d7ff8 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -807,6 +807,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,
@@ -1445,6 +1446,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
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_IF_PRESENT, true))
+ continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
&& gfc_match ("nogroup") == MATCH_YES)
@@ -1955,25 +1961,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
(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_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_SERIAL_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#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_NO_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 \
@@ -2509,7 +2515,7 @@ cleanup:
#define OMP_TASKLOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_SHARED | OMP_CLAUSE_IF \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
| OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_GRAINSIZE \
| OMP_CLAUSE_NUM_TASKS | OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_NOGROUP)
#define OMP_TARGET_CLAUSES \
@@ -2531,7 +2537,7 @@ cleanup:
| OMP_CLAUSE_FROM | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT)
#define OMP_TEAMS_CLAUSES \
(omp_mask (OMP_CLAUSE_NUM_TEAMS) | OMP_CLAUSE_THREAD_LIMIT \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_SHARED | OMP_CLAUSE_REDUCTION)
#define OMP_DISTRIBUTE_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d9dfcabc65e..6c1b1b0aa0e 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2431,6 +2431,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
+ case OMP_MAP_IF_PRESENT:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
+ break;
case OMP_MAP_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 3e470afe32b..700a9352b1b 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11431,6 +11431,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
break;
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
@@ -11841,6 +11842,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
switch (tkind)
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
index 228637f5883..6df5aa65e70 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
@@ -51,6 +51,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v)
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v)
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
index 5d49f6195b8..30c87a91f36 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
@@ -39,6 +39,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
index b94214e8b63..30930a0cf1c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
@@ -111,6 +111,27 @@ contains
!$acc end data
+ !$acc parallel no_create (tip) ! { dg-error "POINTER" }
+ !$acc end parallel
+ !$acc parallel no_create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc end parallel
+ !$acc parallel deviceptr (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copy (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyin (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyout (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end parallel
+ !$acc kernels no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end kernels
+ !$acc data no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end data
+
+
!$acc parallel present (tip) ! { dg-error "POINTER" }
!$acc end parallel
!$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
index f16d62cce69..454417d6a05 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
@@ -7,6 +7,7 @@ program test
logical :: l = .true.
!$acc data if(l) copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end data
@@ -19,7 +20,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index a70f1e737bd..5583ffb4d04 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -8,6 +8,7 @@ program test
!$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end kernels
@@ -25,7 +26,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 2697bb79e7f..e33653bdd78 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -9,6 +9,7 @@ program test
!$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
!$acc end parallel
@@ -28,7 +29,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 1cf7a912133..603617358ae 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -788,6 +788,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_POINTER:
pp_string (pp, "alloc");
break;
+ case GOMP_MAP_IF_PRESENT:
+ pp_string (pp, "no_alloc");
+ break;
case GOMP_MAP_TO:
case GOMP_MAP_TO_PSET:
pp_string (pp, "to");
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 9e356cdfeec..79c5de38db5 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -75,6 +75,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_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_1 | 3),
/* Do not map, copy bits for firstprivate instead. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
diff --git a/libgomp/target.c b/libgomp/target.c
index 84d6daa76ca..467ebc0772b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -667,6 +667,13 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
has_firstprivate = true;
continue;
}
+ else if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
+ {
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = 0;
+ has_firstprivate = true;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -892,6 +899,49 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
continue;
+ case GOMP_MAP_IF_PRESENT:
+ {
+ 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;
+ 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;
+ if (i + 1 < mapnum)
+ {
+ int kind2 = get_kind (short_mapkind, kinds, i + 1);
+ switch (kind2 & typemask)
+ {
+ 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;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ continue;
+ }
default:
break;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
new file mode 100644
index 00000000000..c7a1bd9c015
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
new file mode 100644
index 00000000000..2964a40b217
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c
new file mode 100644
index 00000000000..418ff216612
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c
@@ -0,0 +1,24 @@
+#include <float.h> /* For FLT_EPSILON. */
+#include <math.h> /* For fabs. */
+#include <stdlib.h> /* For abort. */
+
+int main()
+{
+#define N 100
+ float b[N];
+ float c[N];
+
+#pragma acc enter data create(b)
+
+#pragma acc parallel loop no_create(b) no_create(c)
+ for (int i = 0; i < N; ++i)
+ b[i] = i;
+
+#pragma acc exit data copyout(b)
+
+ for (int i = 0; i < N; ++i)
+ if (fabs (b[i] - i) > 10.0*FLT_EPSILON)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
new file mode 100644
index 00000000000..ca9611b777c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
@@ -0,0 +1,35 @@
+! { dg-do run }
+
+! Test no_create clause with data construct when data is present/not present.
+
+program nocreate
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc data no_create (myarr)
+ if (acc_is_present (myarr) .neqv. shared_memory) 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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
new file mode 100644
index 00000000000..16227b8ae22
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
@@ -0,0 +1,71 @@
+! { dg-do run }
+
+! Test no_create clause with data/parallel constructs.
+
+program nocreate
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ call do_on_target(myarr, n)
+
+ do i = 1, n
+ if (shared_memory) then
+ if (myarr(i) .ne. i * 2) stop 1
+ else
+ if (myarr(i) .ne. i) stop 2
+ end if
+ 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 3
+ 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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.f90
new file mode 100644
index 00000000000..739e8356581
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.f90
@@ -0,0 +1,21 @@
+! { dg-do run }
+
+program main
+ implicit none
+ integer i
+ integer, parameter :: n = 100
+ real*4 b(n), c(n)
+ real :: d(n), e(n)
+ common /BLOCK/ d, e
+
+ !$acc enter data create(b) create(d)
+ !$acc parallel loop no_create(b) no_create(c) no_create(/BLOCK/)
+ do i = 1, n
+ b(i) = i
+ d(i) = -i
+ end do
+ !$acc end parallel loop
+ !$acc exit data copyout(b) copyout(d)
+ if (any(abs(b - [(real(i), i = 1, n)]) > 10*epsilon(b))) stop 1
+ if (any(abs(d - [(real(-i), i = 1, n)]) > 10*epsilon(d))) stop 2
+end program main
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-11-15 19:12 ` Tobias Burnus
@ 2019-12-03 15:16 ` Thomas Schwinge
2019-12-03 17:39 ` Tobias Burnus
` (2 more replies)
0 siblings, 3 replies; 12+ messages in thread
From: Thomas Schwinge @ 2019-12-03 15:16 UTC (permalink / raw)
To: Tobias Burnus, Jakub Jelinek; +Cc: gcc-patches, fortran, Julian Brown
[-- Attachment #1.1: Type: text/plain, Size: 8029 bytes --]
Hi!
Jakub, please note question below.
On 2019-11-15T20:11:29+0100, Tobias Burnus <tobias@codesourcery.com> wrote:
> updated version. Changes:
> * Incorporate Thomas's changes
> * Add no_create clause to newly added 'acc serial'
> * Renamed (G)OMP_MAP_NO_ALLOC to (G)OMP_MAP_IF_PRESENT as proposed
> * Make no_create.c effective by adding 'has_firstprivate = true;' to
> target.c.*
Thanks.
> (* If one tries to access c or e in the no_create-3.{c,f90} run-time
> test case, plugin-nvidia rightly complains (illegal memory access),
> using the created 'b' or 'd' works as tested by the test case.
So that's specifically what you fixed above, or is that another problem?
> This
> feature seems to be also broken on the OG9 branch.)
Not surprising, given the insufficient testsuite coverage... ;'-|
I note that you've not addressed the other TODO items that I had put into
the libgomp memory mapping code (see below for reference). I still think
that this should be understood better, that the code as currently
proposed/discussed is "too complex". I have an idea how to do this
differently (easier?), but I still have to sketch that out, and not sure
when I'll get to that. I'm willing to accept that patch as-is, unless
Jakub has any further comments at this point.
Another thing: I've added just another little bit of testsuite coverage,
and another thing broke. See "TODO" in attached incremental patch.
(Please rename the files appropriately.) Please have a look.
This feels like something going wrong in gimplification, when we "Look in
outer OpenACC contexts, to see if there's a data attribute for this
variable" ('gcc/gimplify.c:omp_notice_variable'), but that's just a wild
guess. If you agree/understand that there is a problem, and add some
XFAILed 'gimple' tree-scanning test cases (maybe even just to the libgomp
test cases that I've added), I'm fine to accept that XFAILed, to be
resolved later.
Maybe even that's not specific to the 'no_create' clause, just doesn't
cause any harm (given the existing testsuite...) for other OpenACC
constructs/clauses?
The incremental Fortran test case changes have bene done in a rush; not
sure if they make much sense, or should see some further work applied to
them.
With these items considered/addressed as you feel comfortable, this is OK
for trunk. To record the review effort, please include "Reviewed-by:
Thomas Schwinge <thomas@codesourcery.com>" in the commit log, see
<https://gcc.gnu.org/wiki/Reviewed-by>.
> PS: Remaining bits of the OG9 patch, which are not included are the
> following. I think those are all attach/detach features: a test case for
> "no_create(s.y…)" (i.e. the struct component-ref;
> libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-{3,4}.c) and some
> 'do_detach = false' in libgomp/target.c. Cf. openacc-gcc-9 /…-8 branch
> patch is commit 8e74c2ec2b90819c995444370e742864a685209f of Dec 20,
> 2018. It has been posted as
> https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01418.html
The libgomp memory mapping code:
> 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.
> --- a/include/gomp-constants.h
> +++ b/include/gomp-constants.h
> @@ -75,6 +75,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_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_1 | 3),
> /* Do not map, copy bits for firstprivate instead. */
> GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
> /* Similarly, but store the value in the pointer rather than
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -667,6 +667,13 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> has_firstprivate = true;
> continue;
> }
> + else if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
> + {
> + tgt->list[i].key = NULL;
> + tgt->list[i].offset = 0;
> + has_firstprivate = true;
> + continue;
> + }
> cur_node.host_start = (uintptr_t) hostaddrs[i];
> if (!GOMP_MAP_POINTER_P (kind & typemask))
> cur_node.host_end = cur_node.host_start + sizes[i];
> @@ -892,6 +899,49 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
> + cur_node.host_start - n->host_start;
> continue;
> + case GOMP_MAP_IF_PRESENT:
> + {
> + 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;
> + 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;
> + if (i + 1 < mapnum)
> + {
> + int kind2 = get_kind (short_mapkind, kinds, i + 1);
> + switch (kind2 & typemask)
> + {
> + 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;
> + break;
> + default:
> + break;
> + }
> + }
> + }
> + continue;
> + }
> default:
> break;
> }
My TODO items:
--- libgomp/target.c
+++ libgomp/target.c
@@ -671,6 +671,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
else if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
{
+ //TODO TS is confused. Handling this here, will inhibit 'gomp_map_vars_existing' being used a bit further below.
tgt->list[i].key = NULL;
tgt->list[i].offset = 0;
has_firstprivate = true;
@@ -908,6 +910,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
if (n != NULL)
{
+ //TODO TS is confused. Due to the way the handling of 'GOMP_MAP_NO_ALLOC' is done in the first loop, we're here re-doing 'gomp_map_vars_existing'?
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;
@@ -917,6 +920,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
else
{
+ //TODO This is basically 'GOMP_MAP_FIRSTPRIVATE_INT' handling?
tgt->list[i].key = NULL;
tgt->list[i].offset = OFFSET_INLINED;
tgt->list[i].length = sizes[i];
@@ -928,6 +932,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
switch (kind2 & typemask)
{
case GOMP_MAP_POINTER:
+ //TODO abort();
+ //TODO This code path is exercised by 'libgomp.oacc-fortran/no_create-2.f90'.
+ //TODO TS does not yet understand why this is needed.
+ //TODO Is this somehow similar to 'GOMP_MAP_TO_PSET' handling?
+
/* The data is not present but we have an attach
or pointer clause next. Skip over it. */
i++;
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-Add-OpenACC-2.6-no_create-clause-support-some-.trunk.patch --]
[-- Type: text/x-diff, Size: 11562 bytes --]
From 9a46a8af6374d248c77d6834efaff971da10ecbe Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Mon, 2 Dec 2019 12:53:17 +0100
Subject: [PATCH] Add OpenACC 2.6 `no_create' clause support: some more testing
---
.../libgomp.oacc-c-c++-common/no_create-1.c | 27 ++++--
.../libgomp.oacc-c-c++-common/no_create-1_.c | 82 +++++++++++++++++++
.../libgomp.oacc-c-c++-common/no_create-2.c | 18 ++--
.../libgomp.oacc-c-c++-common/no_create-2_.c | 49 +++++++++++
.../libgomp.oacc-fortran/no_create-1.f90 | 24 +++---
.../libgomp.oacc-fortran/no_create-2.f90 | 47 +++++++----
6 files changed, 206 insertions(+), 41 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1_.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2_.c
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
index c7a1bd9c015..22e0c20cce9 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
@@ -1,4 +1,5 @@
-/* Test no_create clause when data is present on the device. */
+/* Test 'no_create' clause on compute construct, with data present on the
+ device. */
#include <stdlib.h>
#include <stdio.h>
@@ -9,28 +10,36 @@
int
main (int argc, char *argv[])
{
+ int var;
int *arr = (int *) malloc (N * sizeof (*arr));
- int *devptr;
+ int *devptr[2];
+ acc_copyin (&var, sizeof (var));
acc_copyin (arr, N * sizeof (*arr));
- #pragma acc parallel no_create(arr[0:N]) copyout(devptr)
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
{
- devptr = &arr[2];
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
}
-#if !ACC_MEM_SHARED
- if (acc_hostptr (devptr) != (void *) &arr[2])
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
__builtin_abort ();
-#endif
+ acc_delete (&var, sizeof (var));
acc_delete (arr, N * sizeof (*arr));
#if ACC_MEM_SHARED
- if (&arr[2] != devptr)
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
__builtin_abort ();
#else
- if (&arr[2] == devptr)
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
__builtin_abort ();
#endif
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1_.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1_.c
new file mode 100644
index 00000000000..963cb3a68f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1_.c
@@ -0,0 +1,82 @@
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+ acc_copyin (&var, sizeof (var));
+ acc_copyin (arr, N * sizeof (*arr));
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+ }
+
+ acc_delete (&var, sizeof (var));
+ acc_delete (arr, N * sizeof (*arr));
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
index 2964a40b217..fbd01a25956 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
@@ -1,4 +1,5 @@
-/* Test no_create clause when data is not present on the device. */
+/* Test 'no_create' clause on compute construct, with data not present on the
+ device. */
#include <stdlib.h>
#include <stdio.h>
@@ -8,18 +9,19 @@
int
main (int argc, char *argv[])
{
+ int var;
int *arr = (int *) malloc (N * sizeof (*arr));
- int *devptr;
+ int *devptr[2];
- #pragma acc data no_create(arr[0:N])
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
{
- #pragma acc parallel copyout(devptr)
- {
- devptr = &arr[2];
- }
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
}
- if (devptr != &arr[2])
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
__builtin_abort ();
free (arr);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2_.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2_.c
new file mode 100644
index 00000000000..6f0ace501cf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2_.c
@@ -0,0 +1,49 @@
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data not present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+#if ACC_MEM_SHARED
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+#else
+ if (devptr[0] != NULL)
+ __builtin_abort ();
+ if (devptr[1] != NULL)
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } }
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+ }
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
index ca9611b777c..4a1d5da98aa 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
@@ -2,12 +2,12 @@
! Test no_create clause with data construct when data is present/not present.
-program nocreate
+program no_create
use openacc
implicit none
logical :: shared_memory
integer, parameter :: n = 512
- integer :: myarr(n)
+ integer :: myvar, myarr(n)
integer i
shared_memory = .false.
@@ -15,21 +15,25 @@ program nocreate
shared_memory = .true.
!$acc end kernels
+ myvar = 77
do i = 1, n
myarr(i) = 0
end do
- !$acc data no_create (myarr)
- if (acc_is_present (myarr) .neqv. shared_memory) stop 1
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .neqv. shared_memory) stop 10
+ if (acc_is_present (myarr) .neqv. shared_memory) stop 11
!$acc end data
- !$acc enter data copyin (myarr)
- !$acc data no_create (myarr)
- if (acc_is_present (myarr) .eqv. .false.) stop 2
+ !$acc enter data copyin (myvar, myarr)
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .eqv. .false.) stop 20
+ if (acc_is_present (myarr) .eqv. .false.) stop 21
!$acc end data
- !$acc exit data copyout (myarr)
+ !$acc exit data copyout (myvar, myarr)
+ if (myvar .ne. 77) stop 30
do i = 1, n
- if (myarr(i) .ne. 0) stop 3
+ if (myarr(i) .ne. 0) stop 31
end do
-end program nocreate
+end program no_create
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
index 16227b8ae22..0b11f454aca 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
@@ -2,12 +2,12 @@
! Test no_create clause with data/parallel constructs.
-program nocreate
+program no_create
use openacc
implicit none
logical :: shared_memory
integer, parameter :: n = 512
- integer :: myarr(n)
+ integer :: myvar, myarr(n)
integer i
shared_memory = .false.
@@ -15,47 +15,66 @@ program nocreate
shared_memory = .true.
!$acc end kernels
+ myvar = 55
do i = 1, n
myarr(i) = 0
end do
- call do_on_target(myarr, n)
+ call do_on_target(myvar, n, myarr)
+ if (shared_memory) then
+ if (myvar .ne. 44) stop 10
+ else
+ if (myvar .ne. 33) stop 11
+ end if
do i = 1, n
if (shared_memory) then
- if (myarr(i) .ne. i * 2) stop 1
+ if (myarr(i) .ne. i * 2) stop 20
else
- if (myarr(i) .ne. i) stop 2
+ if (myarr(i) .ne. i) stop 21
end if
end do
+ myvar = 55
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)
+ !$acc enter data copyin(myvar, myarr)
+ call do_on_target(myvar, n, myarr)
+ !$acc exit data copyout(myvar, myarr)
+ if (myvar .ne. 44) stop 30
do i = 1, n
- if (myarr(i) .ne. i * 2) stop 3
+ if (myarr(i) .ne. i * 2) stop 31
end do
-end program nocreate
+end program no_create
-subroutine do_on_target (arr, n)
+subroutine do_on_target (var, n, arr)
use openacc
implicit none
- integer :: n, arr(n)
+ integer :: var, n, arr(n)
integer :: i
-!$acc data no_create (arr)
+!$acc data no_create (var, arr)
+if (acc_is_present(var)) 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 !no_create (var)
+ var = 44
+ !$acc end parallel
+else
+ var = 33
+end if
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)
+ !$acc parallel loop !no_create (arr)
do i = 1, n
arr(i) = i * 2
end do
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-12-03 15:16 ` Thomas Schwinge
@ 2019-12-03 17:39 ` Tobias Burnus
2019-12-13 12:07 ` Tobias Burnus
2019-12-17 19:25 ` Tobias Burnus
2 siblings, 0 replies; 12+ messages in thread
From: Tobias Burnus @ 2019-12-03 17:39 UTC (permalink / raw)
To: Thomas Schwinge, Jakub Jelinek; +Cc: gcc-patches, fortran, Julian Brown
On 12/3/19 4:16 PM, Thomas Schwinge wrote:
> On 2019-11-15T20:11:29+0100, Tobias Burnus <tobias@codesourcery.com> wrote:
>> * Make no_create.c effective by adding 'has_firstprivate = true;' to
>> target.c.*
>> (* If one tries to access c or e in the no_create-3.{c,f90} run-time
>> test case, plugin-nvidia rightly complains (illegal memory access),
>> using the created 'b' or 'd' works as tested by the test case.
> So that's specifically what you fixed above, or is that another problem?
Well, that was one way of manually testing that it really worked for
not-mapped variables w/o creating them (i.e. verifying that "no_create"
didn't just act as "present"). – Manual as that's not that simple to
code in the test suite (shared memory, exact wording for dg-output etc.)
— However, I think it can be done using '#include <openacc.h>' / "use
openacc", #if !ACC_MEM_SHARED, and calling acc_is_present (passing
either "sizeof()" or a simple "1" as "len" argument); hence, I will try
this next version of the patch.
> I'm willing to accept that patch as-is, unless Jakub has any further comments at this point. […]
> With these items considered/addressed as you feel comfortable, this is OK for trunk.
Tobias
PS: I will have a closer look tomorrow at the your new test cases and
comments.
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-12-03 15:16 ` Thomas Schwinge
2019-12-03 17:39 ` Tobias Burnus
@ 2019-12-13 12:07 ` Tobias Burnus
2019-12-17 19:25 ` Tobias Burnus
2 siblings, 0 replies; 12+ messages in thread
From: Tobias Burnus @ 2019-12-13 12:07 UTC (permalink / raw)
To: Thomas Schwinge, Tobias Burnus, Harwath, Frederik
Cc: gcc-patches, fortran, Julian Brown
Hi Thomas,
regarding your TODO in your test case about implicit mapping of variables, I did
some testing.
The 'copy' issue is a general feature and not restricted to no_create.
Additionally, 'int *arr' is not a real array: as the compiler does not know the
size, it cannot distinguish a pointer to a scalar integer from a pointer to an
integer array. – OpenACC and OpenMP map 'int *arr' slightly differently.
* * *
Looking at the spec (thanks Frederik for the help), I read it such that
* For OpenACC 2.6+2.7 in both kernels (2.5.2) and parallel constructs (2.5.1)
[for both, see last paragraph of 'Description']:
– with 'default(none)': nothing is done explicitly.
– with 'default(present)': then scalars = 'copy', arrays/combined types 'present'
- otherwise: arrays/combined types = 'copy' and
parallel: scalars = 'firstprivate'
kernels: scalars = 'copy'
(Per definition, Fortran's allocatable, pointer + character are never a 'scalar'.)
 * For OpenMP, implicit mappings is handled similar to 'parallel':
– scalars = firstprivate (unless: 'defaultmap(tofrom:scalars)')
– nonscalars = 'map(tofrom:'
(OpenMP 5 permits more 'defaultmap's and Fortran allocatable/pointer scalars are
then also 'map(tofrom:' by default; note Fortran's 'character' is not a 'scalar'
per OpenMP terminology.)
For 'int *arr', one has a pointer which can point to a single or multiple
("array") integer – the in C/C++ compiler cannot know, contrary to 'int arr2[4]'.
Assume now: 'int var, *arr, arr2[4]' (all -fdump-tree-omplower). Result:
(A) OpenACC
oacc_parallel map(tofrom:arr2 [len: 16]) firstprivate(arr) firstprivate(var)
oacc_kernels map(tofrom:arr2 [len: 16]) map(force_tofrom:arr [len: 8]) map(force_tofrom:var [len: 4])
(B) OpenMP
omp target map(tofrom:arr2 [len: 16]) \
map(alloc:MEM[(char *)arr] [len: 0]) map(firstprivate:arr [pointer assign, bias: 0]) \
firstprivate(var)
Which looks fine – despite the difference between OpenMP and OpenACC.
(OpenACC: Using default(present) also works – giving 'map(force_present:arr2'; as does
default(none) – causing the compiler to complain about unmapped variables.)
* * *
When enclosing this in 'acc data' (or 'omp data target'), the following of OpenACC applies:
'implicitly determine data attributes for variables that are referenced
in the compute construct that […] do not appear in a data clause on […]
a lexically containing data construct […]".
Testing shows that independent of the used clause, 'copy()' is always done, also
for scalars in 'parallel'.
For OpenMP 4.5, 2.15.5 is a bit unclear whether 'omp data target's map() apply or not,
but GCC currently ignores them completely and does the normal 'map(tofrom:' + 'firstprivate'
mapping in this case.
Tobias
PS: Your example was:
On 12/3/19 4:16 PM, Thomas Schwinge wrote:
> + int var;
> + int *arr = (int *) malloc (N * sizeof (*arr));
> + int *devptr[2];
> +
> +#pragma acc data no_create(var, arr[0:N])
> + {
> + devptr[0] = (int *) acc_deviceptr (&var);
> + devptr[1] = (int *) acc_deviceptr (&arr[2]);
> +
> +#if ACC_MEM_SHARED
> + if (devptr[0] == NULL)
> + __builtin_abort ();
> + if (devptr[1] == NULL)
> + __builtin_abort ();
> +#else
> + if (devptr[0] != NULL)
> + __builtin_abort ();
> + if (devptr[1] != NULL)
> + __builtin_abort ();
> +#endif
> +
> +#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
> + {
> + devptr[0] = &var;
> + devptr[1] = &arr[2];
> + }
> +
> + if (devptr[0] != &var)
> + __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } }
> + if (devptr[1] != &arr[2])
> + __builtin_abort ();
> + }
> +
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-12-03 15:16 ` Thomas Schwinge
2019-12-03 17:39 ` Tobias Burnus
2019-12-13 12:07 ` Tobias Burnus
@ 2019-12-17 19:25 ` Tobias Burnus
2019-12-18 12:41 ` Tobias Burnus
2 siblings, 1 reply; 12+ messages in thread
From: Tobias Burnus @ 2019-12-17 19:25 UTC (permalink / raw)
To: Thomas Schwinge, Jakub Jelinek; +Cc: gcc-patches, fortran, Julian Brown
[-- Attachment #1: Type: text/plain, Size: 4772 bytes --]
Hi Thomas,
I am reasonably comfortable with the current patch (regarding your
TODOs) – see attachment. It is the previous patch plus your changes plus
one additional condition (see below) in target.c's first
GOMP_MAP_IF_PRESENT handling.
I intent to re-test it tomorrow and then commit it, unless some other
issues or comments come up. — See a bunch of comments below.
Cheers,
Tobias
On 12/3/19 4:16 PM, Thomas Schwinge wrote:
> So that's specifically what you fixed above
(See previous reply in this email. Now added an acc_is_present check.
https://gcc.gnu.org/ml/gcc-patches/2019-12/msg00156.html)
> Another thing: I've added just another little bit of testsuite
> coverage, and another thing broke. See "TODO" in attached incremental
> patch. […]
Files included, the other issue was XFAILed by you (and hence passed). A
fix for that issue is:
https://gcc.gnu.org/ml/gcc-patches/2019-12/msg01135.html — and a
completely separate issue. (That patch is small, very localized and
orthogonal to this patch.)
> The incremental Fortran test case changes have bene done in a rush; not
> sure if they make much sense, or should see some further work applied to
> them.
I think one can do more, but they are fine. I am not 100% sure how to
read the following:
 ! 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 !no_create (var)
First, why is 'no_create(var)' now commented? – For this code, it should
really work both ways and independent whether commented boils down to
'copy' (currently) or 'present' (with my other patch, linked above).
> With these items considered/addressed as you feel comfortable, this is OK
> for trunk.
> My TODO items:
>
> --- libgomp/target.c
> +++ libgomp/target.c
> @@ -671,6 +671,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> }
> else if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
> {
> + //TODO TS is confused. Handling this here, will inhibit 'gomp_map_vars_existing' being used a bit further below.
> tgt->list[i].key = NULL;
> tgt->list[i].offset = 0;
> has_firstprivate = true;
True – but should it? the only effect seems to be that it bumps the ref
count. (Should it or shouldn't it?) In any case if the data is not
present, it will fail in this section.
However, I think the following is missing before 'continue' – even
though testing did not hit it:
/* Handle the attach/pointer clause next to it later, together with
GOMP_MAP_IF_PRESENT as the data might be not available. */
if (i + 1 < mapnum
&& ((typemask & get_kind (short_mapkind, kinds, i + 1))
== GOMP_MAP_POINTER))
++i;
> @@ -908,6 +910,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
> if (n != NULL)
> {
> + //TODO TS is confused. Due to the way the handling of 'GOMP_MAP_NO_ALLOC' is done in the first loop, we're here re-doing 'gomp_map_vars_existing'?
> 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;
Essentially, yes – except that we know here that the variable does exist
– in the block above, it also works, but only if the variable has been
mapped at some point.
> @@ -917,6 +920,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> }
> else
> {
> + //TODO This is basically 'GOMP_MAP_FIRSTPRIVATE_INT' handling?
> tgt->list[i].key = NULL;
> tgt->list[i].offset = OFFSET_INLINED;
> tgt->list[i].length = sizes[i];
Yes – but one could also call it 'hostaddrs[i] == NULL' handling, which
makes more sense semantically.
> @@ -928,6 +932,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> switch (kind2 & typemask)
> {
> case GOMP_MAP_POINTER:
> + //TODO abort();
> + //TODO This code path is exercised by 'libgomp.oacc-fortran/no_create-2.f90'.
> + //TODO TS does not yet understand why this is needed.
> + //TODO Is this somehow similar to 'GOMP_MAP_TO_PSET' handling?
> +
> /* The data is not present but we have an attach
> or pointer clause next. Skip over it. */
> i++;
Yes, as -fdump-tree-omplower shows, it is handled like a normal map,
except that the variable itself gets a 'no_alloc'.
map(no_alloc:*var.7_5 [len: 4]) map(alloc:var [pointer assign, bias: 0])
map(no_alloc:(*arr.8_6)
[-- Attachment #2: openacc_no_create8.diff --]
[-- Type: text/x-patch, Size: 39154 bytes --]
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.
2019-12-18 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@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): 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_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
* trans-openmp.c (gfc_trans_omp_clauses_1): Support
OMP_MAP_NO_ALLOC.
gcc/testsuite/
* gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
* gfortran.dg/goacc/common-block-1.f90: Likewise.
* gfortran.dg/goacc/data-clauses.f95: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
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/no_create-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test.
* testsuite/libgomp.oacc-fortran/no_create-1.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-2.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-3.f90: New test.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
gcc/c-family/c-pragma.h | 1 +
gcc/c/c-parser.c | 20 ++++-
gcc/c/c-typeck.c | 1 +
gcc/cp/parser.c | 22 +++++-
gcc/cp/semantics.c | 1 +
gcc/fortran/gfortran.h | 1 +
gcc/fortran/openmp.c | 28 ++++---
gcc/fortran/trans-openmp.c | 3 +
gcc/omp-low.c | 2 +
gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 | 3 +
gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 | 3 +
gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 | 21 +++++
gcc/testsuite/gfortran.dg/goacc/data-tree.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 | 3 +-
gcc/tree-pretty-print.c | 3 +
include/gomp-constants.h | 2 +
libgomp/target.c | 56 ++++++++++++
.../libgomp.oacc-c-c++-common/no_create-1.c | 49 ++++++++++++
.../libgomp.oacc-c-c++-common/no_create-2.c | 30 ++++++++
.../libgomp.oacc-c-c++-common/no_create-3.c | 37 +++++++++
.../libgomp.oacc-c-c++-common/no_create-4.c | 82 ++++++++++++++++++++
.../libgomp.oacc-c-c++-common/no_create-5.c | 49 ++++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-1.f90 | 39 ++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-2.f90 | 90 ++++++++++++++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-3.f90 | 34 ++++++++
26 files changed, 561 insertions(+), 19 deletions(-)
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index bfe681bb430..3754c5fda45 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -154,6 +154,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_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index bfe56998996..9b8008816d2 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12650,7 +12650,9 @@ c_parser_omp_clause_name (c_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -13113,7 +13115,10 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -13149,6 +13154,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -15947,6 +15955,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
clauses = c_parser_oacc_single_int_clause (parser,
OMP_CLAUSE_NUM_GANGS,
@@ -16415,6 +16427,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
| (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
@@ -16747,6 +16760,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -16762,6 +16776,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -16780,6 +16795,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 36aedc063d2..ce5e6495fb1 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13422,6 +13422,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 05be440cb9b..d024ea76366 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -33191,7 +33191,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -33557,7 +33559,10 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -33593,6 +33598,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -36155,6 +36163,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@@ -39960,6 +39972,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
| (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
@@ -40281,6 +40294,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -40295,8 +40309,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
+ | (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) \
@@ -40314,6 +40329,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 83a7891e57b..8b2e258e4bb 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5288,6 +5288,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index f4a2b99bdc4..3907d1407ac 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1192,6 +1192,7 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_IF_PRESENT,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index dc0521b40f0..576003d7ff8 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -807,6 +807,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,
@@ -1445,6 +1446,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
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_IF_PRESENT, true))
+ continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
&& gfc_match ("nogroup") == MATCH_YES)
@@ -1955,25 +1961,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
(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_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_SERIAL_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#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_NO_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 \
@@ -2509,7 +2515,7 @@ cleanup:
#define OMP_TASKLOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_SHARED | OMP_CLAUSE_IF \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
| OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_GRAINSIZE \
| OMP_CLAUSE_NUM_TASKS | OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_NOGROUP)
#define OMP_TARGET_CLAUSES \
@@ -2531,7 +2537,7 @@ cleanup:
| OMP_CLAUSE_FROM | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT)
#define OMP_TEAMS_CLAUSES \
(omp_mask (OMP_CLAUSE_NUM_TEAMS) | OMP_CLAUSE_THREAD_LIMIT \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_SHARED | OMP_CLAUSE_REDUCTION)
#define OMP_DISTRIBUTE_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 0649a34b9eb..0053257f2cb 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2624,6 +2624,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
+ case OMP_MAP_IF_PRESENT:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
+ break;
case OMP_MAP_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ad26f7918a5..08ce00603ca 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11431,6 +11431,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
break;
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
@@ -11842,6 +11843,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
switch (tkind)
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
index 228637f5883..6df5aa65e70 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
@@ -51,6 +51,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v)
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v)
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
index 5d49f6195b8..30c87a91f36 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
@@ -39,6 +39,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
index b94214e8b63..30930a0cf1c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
@@ -111,6 +111,27 @@ contains
!$acc end data
+ !$acc parallel no_create (tip) ! { dg-error "POINTER" }
+ !$acc end parallel
+ !$acc parallel no_create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc end parallel
+ !$acc parallel deviceptr (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copy (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyin (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyout (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end parallel
+ !$acc kernels no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end kernels
+ !$acc data no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end data
+
+
!$acc parallel present (tip) ! { dg-error "POINTER" }
!$acc end parallel
!$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
index f16d62cce69..454417d6a05 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
@@ -7,6 +7,7 @@ program test
logical :: l = .true.
!$acc data if(l) copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end data
@@ -19,7 +20,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index a70f1e737bd..5583ffb4d04 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -8,6 +8,7 @@ program test
!$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end kernels
@@ -25,7 +26,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 2697bb79e7f..e33653bdd78 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -9,6 +9,7 @@ program test
!$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
!$acc end parallel
@@ -28,7 +29,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 1cf7a912133..603617358ae 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -788,6 +788,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_POINTER:
pp_string (pp, "alloc");
break;
+ case GOMP_MAP_IF_PRESENT:
+ pp_string (pp, "no_alloc");
+ break;
case GOMP_MAP_TO:
case GOMP_MAP_TO_PSET:
pp_string (pp, "to");
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 9e356cdfeec..79c5de38db5 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -75,6 +75,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_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_1 | 3),
/* Do not map, copy bits for firstprivate instead. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
diff --git a/libgomp/target.c b/libgomp/target.c
index 84d6daa76ca..467ebc0772b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -667,6 +667,19 @@ gomp_map_vars_internal (struct gomp_device_descr *
has_firstprivate = true;
continue;
}
+ else if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
+ {
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = 0;
+ has_firstprivate = true;
+ /* Handle the attach/pointer clause next to it later, together with
+ GOMP_MAP_IF_PRESENT as the data might be not available. */
+ if (i + 1 < mapnum
+ && ((typemask & get_kind (short_mapkind, kinds, i + 1))
+ == GOMP_MAP_POINTER))
+ ++i;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -892,6 +905,49 @@ gomp_map_vars_internal (struct gomp_device_descr *
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
continue;
+ case GOMP_MAP_IF_PRESENT:
+ {
+ 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;
+ 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;
+ if (i + 1 < mapnum)
+ {
+ int kind2 = get_kind (short_mapkind, kinds, i + 1);
+ switch (kind2 & typemask)
+ {
+ 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;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ continue;
+ }
default:
break;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
new file mode 100644
index 00000000000..22e0c20cce9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
@@ -0,0 +1,49 @@
+/* Test 'no_create' clause on compute construct, with data present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+ acc_copyin (&var, sizeof (var));
+ acc_copyin (arr, N * sizeof (*arr));
+
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+ acc_delete (&var, sizeof (var));
+ acc_delete (arr, N * sizeof (*arr));
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
new file mode 100644
index 00000000000..fbd01a25956
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
@@ -0,0 +1,30 @@
+/* Test 'no_create' clause on compute construct, with data not present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c
new file mode 100644
index 00000000000..d21f1d18600
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c
@@ -0,0 +1,37 @@
+#include <float.h> /* For FLT_EPSILON. */
+#include <math.h> /* For fabs. */
+#include <stdlib.h> /* For abort. */
+
+#include <openacc.h> /* For acc_is_present. */
+
+
+int main()
+{
+#define N 100
+ float b[N];
+ float c[N];
+
+#pragma acc enter data create(b)
+
+ if (!acc_is_present(b))
+ abort();
+ if (acc_is_present(c))
+ abort();
+
+#pragma acc parallel loop no_create(b) no_create(c)
+ for (int i = 0; i < N; ++i)
+ b[i] = i;
+
+ if (!acc_is_present(b))
+ abort();
+ if (acc_is_present(c))
+ abort();
+
+#pragma acc exit data copyout(b)
+
+ for (int i = 0; i < N; ++i)
+ if (fabs (b[i] - i) > 10.0*FLT_EPSILON)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c
new file mode 100644
index 00000000000..963cb3a68f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c
@@ -0,0 +1,82 @@
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+ acc_copyin (&var, sizeof (var));
+ acc_copyin (arr, N * sizeof (*arr));
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+ }
+
+ acc_delete (&var, sizeof (var));
+ acc_delete (arr, N * sizeof (*arr));
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c
new file mode 100644
index 00000000000..6f0ace501cf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c
@@ -0,0 +1,49 @@
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data not present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+#if ACC_MEM_SHARED
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+#else
+ if (devptr[0] != NULL)
+ __builtin_abort ();
+ if (devptr[1] != NULL)
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } }
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+ }
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
new file mode 100644
index 00000000000..4a1d5da98aa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
@@ -0,0 +1,39 @@
+! { dg-do run }
+
+! Test no_create clause with data construct when data is present/not present.
+
+program no_create
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myvar, myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ myvar = 77
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .neqv. shared_memory) stop 10
+ if (acc_is_present (myarr) .neqv. shared_memory) stop 11
+ !$acc end data
+
+ !$acc enter data copyin (myvar, myarr)
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .eqv. .false.) stop 20
+ if (acc_is_present (myarr) .eqv. .false.) stop 21
+ !$acc end data
+ !$acc exit data copyout (myvar, myarr)
+
+ if (myvar .ne. 77) stop 30
+ do i = 1, n
+ if (myarr(i) .ne. 0) stop 31
+ end do
+end program no_create
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
new file mode 100644
index 00000000000..0b11f454aca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
@@ -0,0 +1,90 @@
+! { dg-do run }
+
+! Test no_create clause with data/parallel constructs.
+
+program no_create
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myvar, myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ myvar = 55
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ call do_on_target(myvar, n, myarr)
+
+ if (shared_memory) then
+ if (myvar .ne. 44) stop 10
+ else
+ if (myvar .ne. 33) stop 11
+ end if
+ do i = 1, n
+ if (shared_memory) then
+ if (myarr(i) .ne. i * 2) stop 20
+ else
+ if (myarr(i) .ne. i) stop 21
+ end if
+ end do
+
+ myvar = 55
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc enter data copyin(myvar, myarr)
+ call do_on_target(myvar, n, myarr)
+ !$acc exit data copyout(myvar, myarr)
+
+ if (myvar .ne. 44) stop 30
+ do i = 1, n
+ if (myarr(i) .ne. i * 2) stop 31
+ end do
+end program no_create
+
+subroutine do_on_target (var, n, arr)
+ use openacc
+ implicit none
+ integer :: var, n, arr(n)
+ integer :: i
+
+!$acc data no_create (var, arr)
+
+if (acc_is_present(var)) 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 !no_create (var)
+ var = 44
+ !$acc end parallel
+else
+ var = 33
+end if
+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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.f90
new file mode 100644
index 00000000000..f69e43ca998
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+
+program main
+ use openacc, only: acc_is_present
+ implicit none
+ integer i
+ integer, parameter :: n = 100
+ real*4 b(n), c(n)
+ real :: d(n), e(n)
+ common /BLOCK/ d, e
+
+ !$acc enter data create(b) create(d)
+
+ if (.not. acc_is_present(b)) stop 1
+ if (acc_is_present(c)) stop 2
+ if (.not. acc_is_present(d)) stop 3
+ if (acc_is_present(e)) stop 4
+
+ !$acc parallel loop no_create(b) no_create(c) no_create(/BLOCK/)
+ do i = 1, n
+ b(i) = i
+ d(i) = -i
+ end do
+ !$acc end parallel loop
+
+ if (.not. acc_is_present(b)) stop 5
+ if (acc_is_present(c)) stop 6
+ if (.not. acc_is_present(d)) stop 7
+ if (acc_is_present(e)) stop 8
+
+ !$acc exit data copyout(b) copyout(d)
+ if (any(abs(b - [(real(i), i = 1, n)]) > 10*epsilon(b))) stop 9
+ if (any(abs(d - [(real(-i), i = 1, n)]) > 10*epsilon(d))) stop 10
+end program main
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-12-17 19:25 ` Tobias Burnus
@ 2019-12-18 12:41 ` Tobias Burnus
2019-12-18 21:47 ` Thomas Schwinge
0 siblings, 1 reply; 12+ messages in thread
From: Tobias Burnus @ 2019-12-18 12:41 UTC (permalink / raw)
To: Thomas Schwinge, Jakub Jelinek; +Cc: gcc-patches, fortran, Julian Brown
[-- Attachment #1: Type: text/plain, Size: 7919 bytes --]
Hi Thomas,
@Thomas (and, possibly, Julian & Jakub): Please glance quickly the
gomp_map_vars_internal change.
libgomp/target.c's gomp_map_vars_internal: it now uses the normal code
path in the upper loop, except that one directly bails out when the
'key' has not been found (skipping the adjacent MAP_POINTER as well).
The 'case' in the second loop is only reached, if tgt[i]->key == NULL
(i.e. if not present) and one can unconditionally skip here. — This
seems to be cleaner and should avoid some confusions :-)
GOMP_MAP_POINTER, following MAP_IF_PRESENT: I am not sure about this.
The testsuite digests both mapping and skipping the map pointer. It
looks a tad cleaner to avoid mapping the pointer (if the var is not
present) – saving also few bytes and cpu cycles. On the down side, it
adds an order dependence assumption, namely assuming that the
MAP_POINTER after 'no_create'/MAP_IF_PRESENT always belongs to
no_create. – [This patch follows the original patch and skips the
map_pointer.]
Otherwise, except for added acc_is_present calls to no_create-3.c to
check that no_create does not cause mapping and applying your/Thomas's
patches, it matches my previous version, which was OK'ed. — Hence, I
intent to commit it tomorrow, unless there are further comments.
Cheers,
Tobias
On 12/17/19 8:11 PM, Tobias Burnus wrote:
> Hi Thomas,
>
> I am reasonably comfortable with the current patch (regarding your
> TODOs) – see attachment. It is the previous patch plus your changes
> plus one additional condition (see below) in target.c's first
> GOMP_MAP_IF_PRESENT handling.
>
> I intent to re-test it tomorrow and then commit it, unless some other
> issues or comments come up. — See a bunch of comments below.
>
> Cheers,
>
> Tobias
>
> On 12/3/19 4:16 PM, Thomas Schwinge wrote:
>> So that's specifically what you fixed above
> (See previous reply in this email. Now added an acc_is_present check.
> https://gcc.gnu.org/ml/gcc-patches/2019-12/msg00156.html)
>> Another thing: I've added just another little bit of testsuite
>> coverage, and another thing broke. See "TODO" in attached incremental
>> patch. […]
> Files included, the other issue was XFAILed by you (and hence passed).
> A fix for that issue is:
> https://gcc.gnu.org/ml/gcc-patches/2019-12/msg01135.html — and a
> completely separate issue. (That patch is small, very localized and
> orthogonal to this patch.)
>> The incremental Fortran test case changes have bene done in a rush; not
>> sure if they make much sense, or should see some further work applied to
>> them.
>
> I think one can do more, but they are fine. I am not 100% sure how to
> read the following:
>
> Â ! 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 !no_create (var)
>
> First, why is 'no_create(var)' now commented? – For this code, it
> should really work both ways and independent whether commented boils
> down to 'copy' (currently) or 'present' (with my other patch, linked
> above).
>
>> With these items considered/addressed as you feel comfortable, this
>> is OK
>> for trunk.
>
>> My TODO items:
>>
>> --- libgomp/target.c
>> +++ libgomp/target.c
>> @@ -671,6 +671,7 @@ gomp_map_vars_internal (struct gomp_device_descr
>> *devicep,
>> Â Â Â Â Â }
>> Â Â Â Â Â Â Â else if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
>> Â Â Â Â Â {
>> +     //TODO TS is confused. Handling this here, will inhibit
>> 'gomp_map_vars_existing' being used a bit further below.
>> Â Â Â Â Â Â Â tgt->list[i].key = NULL;
>> Â Â Â Â Â Â Â tgt->list[i].offset = 0;
>> Â Â Â Â Â Â Â has_firstprivate = true;
>
> True – but should it? the only effect seems to be that it bumps the
> ref count. (Should it or shouldn't it?) In any case if the data is not
> present, it will fail in this section.
>
> However, I think the following is missing before 'continue' – even
> though testing did not hit it:
>
> Â Â Â Â Â /* Handle the attach/pointer clause next to it later, together with
> Â Â Â Â Â Â Â Â GOMP_MAP_IF_PRESENT as the data might be not available. */
> Â Â Â Â Â if (i + 1 < mapnum
> Â Â Â Â Â Â Â Â Â && ((typemask & get_kind (short_mapkind, kinds, i + 1))
> Â Â Â Â Â Â Â Â Â == GOMP_MAP_POINTER))
> Â Â Â Â Â Â Â ++i;
>
>> @@ -908,6 +910,7 @@ gomp_map_vars_internal (struct gomp_device_descr
>> *devicep,
>> Â Â Â Â Â Â Â Â Â Â Â splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
>> Â Â Â Â Â Â Â Â Â Â Â if (n != NULL)
>> Â Â Â Â Â Â Â Â Â Â Â Â Â {
>> +             //TODO TS is confused. Due to the way the handling of
>> 'GOMP_MAP_NO_ALLOC' is done in the first loop, we're here re-doing
>> 'gomp_map_vars_existing'?
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â 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;
> Essentially, yes – except that we know here that the variable does
> exist – in the block above, it also works, but only if the variable
> has been mapped at some point.
>> @@ -917,6 +920,7 @@ gomp_map_vars_internal (struct gomp_device_descr
>> *devicep,
>> Â Â Â Â Â Â Â Â Â Â Â Â Â }
>> Â Â Â Â Â Â Â Â Â Â Â else
>> Â Â Â Â Â Â Â Â Â Â Â Â Â {
>> +Â Â Â Â Â Â Â Â Â Â Â Â Â //TODO This is basically 'GOMP_MAP_FIRSTPRIVATE_INT'
>> handling?
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â tgt->list[i].key = NULL;
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â tgt->list[i].offset = OFFSET_INLINED;
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â tgt->list[i].length = sizes[i];
> Yes – but one could also call it 'hostaddrs[i] == NULL' handling,
> which makes more sense semantically.
>> @@ -928,6 +932,11 @@ gomp_map_vars_internal (struct gomp_device_descr
>> *devicep,
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â switch (kind2 & typemask)
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â {
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â case GOMP_MAP_POINTER:
>> +Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â //TODO abort();
>> +Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â //TODO This code path is exercised by
>> 'libgomp.oacc-fortran/no_create-2.f90'.
>> +Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â //TODO TS does not yet understand why this is needed.
>> +Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â //TODO Is this somehow similar to
>> 'GOMP_MAP_TO_PSET' handling?
>> +
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â /* The data is not present but we have an attach
>>                   or pointer clause next. Skip over it. */
>> Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â Â i++;
>
> Yes, as -fdump-tree-omplower shows, it is handled like a normal map,
> except that the variable itself gets a 'no_alloc'.
>
> map(no_alloc:*var.7_5 [len: 4]) map(alloc:var [pointer assign, bias:
> 0]) map(no_alloc:(*arr.8_6)
>
[-- Attachment #2: openacc_no_create9.diff --]
[-- Type: text/x-patch, Size: 38058 bytes --]
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.
2019-12-18 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@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): 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_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
* trans-openmp.c (gfc_trans_omp_clauses_1): Support
OMP_MAP_NO_ALLOC.
gcc/testsuite/
* gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
* gfortran.dg/goacc/common-block-1.f90: Likewise.
* gfortran.dg/goacc/data-clauses.f95: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
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/no_create-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test.
* testsuite/libgomp.oacc-fortran/no_create-1.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-2.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-3.F90: New test.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
gcc/c-family/c-pragma.h | 1 +
gcc/c/c-parser.c | 20 ++++-
gcc/c/c-typeck.c | 1 +
gcc/cp/parser.c | 22 +++++-
gcc/cp/semantics.c | 1 +
gcc/fortran/gfortran.h | 1 +
gcc/fortran/openmp.c | 28 ++++---
gcc/fortran/trans-openmp.c | 3 +
gcc/omp-low.c | 2 +
gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 | 3 +
gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 | 3 +
gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 | 21 +++++
gcc/testsuite/gfortran.dg/goacc/data-tree.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 | 3 +-
gcc/tree-pretty-print.c | 3 +
include/gomp-constants.h | 2 +
libgomp/target.c | 23 ++++++
.../libgomp.oacc-c-c++-common/no_create-1.c | 49 ++++++++++++
.../libgomp.oacc-c-c++-common/no_create-2.c | 30 ++++++++
.../libgomp.oacc-c-c++-common/no_create-3.c | 25 ++++++
.../libgomp.oacc-c-c++-common/no_create-4.c | 82 ++++++++++++++++++++
.../libgomp.oacc-c-c++-common/no_create-5.c | 49 ++++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-1.f90 | 39 ++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-2.f90 | 90 ++++++++++++++++++++++
.../testsuite/libgomp.oacc-fortran/no_create-3.F90 | 39 ++++++++++
26 files changed, 527 insertions(+), 19 deletions(-)
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index bfe681bb430..3754c5fda45 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -154,6 +154,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_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index bfe56998996..9b8008816d2 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12650,7 +12650,9 @@ c_parser_omp_clause_name (c_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -13113,7 +13115,10 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -13149,6 +13154,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -15947,6 +15955,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
clauses = c_parser_oacc_single_int_clause (parser,
OMP_CLAUSE_NUM_GANGS,
@@ -16415,6 +16427,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
| (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
@@ -16747,6 +16760,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -16762,6 +16776,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -16780,6 +16795,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 36aedc063d2..ce5e6495fb1 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13422,6 +13422,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index f61089934df..c68e01800d4 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -33617,7 +33617,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
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 ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@@ -33983,7 +33985,10 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
copyout ( variable-list )
create ( variable-list )
delete ( 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,
@@ -34019,6 +34024,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -36581,6 +36589,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
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_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@@ -40386,6 +40398,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
| (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
@@ -40707,6 +40720,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
@@ -40721,8 +40735,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
+ | (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) \
@@ -40740,6 +40755,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (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) \
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 42611682549..d6aa94eca04 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5288,6 +5288,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index f4a2b99bdc4..3907d1407ac 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1192,6 +1192,7 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_IF_PRESENT,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index dc0521b40f0..576003d7ff8 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -807,6 +807,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,
@@ -1445,6 +1446,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
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_IF_PRESENT, true))
+ continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
&& gfc_match ("nogroup") == MATCH_YES)
@@ -1955,25 +1961,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
(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_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_SERIAL_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#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_NO_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 \
@@ -2509,7 +2515,7 @@ cleanup:
#define OMP_TASKLOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_SHARED | OMP_CLAUSE_IF \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
| OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_GRAINSIZE \
| OMP_CLAUSE_NUM_TASKS | OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_NOGROUP)
#define OMP_TARGET_CLAUSES \
@@ -2531,7 +2537,7 @@ cleanup:
| OMP_CLAUSE_FROM | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT)
#define OMP_TEAMS_CLAUSES \
(omp_mask (OMP_CLAUSE_NUM_TEAMS) | OMP_CLAUSE_THREAD_LIMIT \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_SHARED | OMP_CLAUSE_REDUCTION)
#define OMP_DISTRIBUTE_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index b6da7b983d5..7153491a460 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2624,6 +2624,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
+ case OMP_MAP_IF_PRESENT:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
+ break;
case OMP_MAP_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index d422c205836..deed83b8c33 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11431,6 +11431,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
break;
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
@@ -11842,6 +11843,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
switch (tkind)
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
index 228637f5883..6df5aa65e70 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
@@ -51,6 +51,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v)
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v)
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
index 5d49f6195b8..30c87a91f36 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
@@ -39,6 +39,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end parallel
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
index b94214e8b63..30930a0cf1c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
@@ -111,6 +111,27 @@ contains
!$acc end data
+ !$acc parallel no_create (tip) ! { dg-error "POINTER" }
+ !$acc end parallel
+ !$acc parallel no_create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc end parallel
+ !$acc parallel deviceptr (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copy (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyin (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyout (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end parallel
+ !$acc kernels no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end kernels
+ !$acc data no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end data
+
+
!$acc parallel present (tip) ! { dg-error "POINTER" }
!$acc end parallel
!$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
index f16d62cce69..454417d6a05 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
@@ -7,6 +7,7 @@ program test
logical :: l = .true.
!$acc data if(l) copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end data
@@ -19,7 +20,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index a70f1e737bd..5583ffb4d04 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -8,6 +8,7 @@ program test
!$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end kernels
@@ -25,7 +26,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 2697bb79e7f..e33653bdd78 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -9,6 +9,7 @@ program test
!$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
!$acc end parallel
@@ -28,7 +29,7 @@ end program test
! { 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\\(no_alloc:n\\)" 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" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 1cf7a912133..603617358ae 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -788,6 +788,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_POINTER:
pp_string (pp, "alloc");
break;
+ case GOMP_MAP_IF_PRESENT:
+ pp_string (pp, "no_alloc");
+ break;
case GOMP_MAP_TO:
case GOMP_MAP_TO_PSET:
pp_string (pp, "to");
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 9e356cdfeec..79c5de38db5 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -75,6 +75,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_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_1 | 3),
/* Do not map, copy bits for firstprivate instead. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
diff --git a/libgomp/target.c b/libgomp/target.c
index 82ed38c01ec..9febd0ebc15 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -706,6 +706,21 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
{
tgt->list[i].key = NULL;
+ if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
+ {
+ /* Not present, hence, skip entry - including its MAP_POINTER,
+ when existing. */
+ tgt->list[i].offset = 0;
+ if (i + 1 < mapnum
+ && ((typemask & get_kind (short_mapkind, kinds, i + 1))
+ == GOMP_MAP_POINTER))
+ {
+ ++i;
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = 0;
+ }
+ continue;
+ }
size_t align = (size_t) 1 << (kind >> rshift);
not_found_cnt++;
if (tgt_align < align)
@@ -892,6 +907,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
continue;
+ case GOMP_MAP_IF_PRESENT:
+ /* Not present - otherwise handled above. Skip over its
+ MAP_POINTER as well. */
+ if (i + 1 < mapnum
+ && ((typemask & get_kind (short_mapkind, kinds, i + 1))
+ == GOMP_MAP_POINTER))
+ ++i;
+ continue;
default:
break;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
new file mode 100644
index 00000000000..22e0c20cce9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c
@@ -0,0 +1,49 @@
+/* Test 'no_create' clause on compute construct, with data present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+ acc_copyin (&var, sizeof (var));
+ acc_copyin (arr, N * sizeof (*arr));
+
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+ acc_delete (&var, sizeof (var));
+ acc_delete (arr, N * sizeof (*arr));
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
new file mode 100644
index 00000000000..fbd01a25956
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c
@@ -0,0 +1,30 @@
+/* Test 'no_create' clause on compute construct, with data not present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c
new file mode 100644
index 00000000000..18466b88b5c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c
@@ -0,0 +1,25 @@
+#include <float.h> /* For FLT_EPSILON. */
+#include <math.h> /* For fabs. */
+#include <stdlib.h> /* For abort. */
+
+
+int main()
+{
+#define N 100
+ float b[N];
+ float c[N];
+
+#pragma acc enter data create(b)
+
+#pragma acc parallel loop no_create(b) no_create(c)
+ for (int i = 0; i < N; ++i)
+ b[i] = i;
+
+#pragma acc exit data copyout(b)
+
+ for (int i = 0; i < N; ++i)
+ if (fabs (b[i] - i) > 10.0*FLT_EPSILON)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c
new file mode 100644
index 00000000000..963cb3a68f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c
@@ -0,0 +1,82 @@
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+ acc_copyin (&var, sizeof (var));
+ acc_copyin (arr, N * sizeof (*arr));
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+ }
+
+ acc_delete (&var, sizeof (var));
+ acc_delete (arr, N * sizeof (*arr));
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c
new file mode 100644
index 00000000000..6f0ace501cf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c
@@ -0,0 +1,49 @@
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data not present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+#if ACC_MEM_SHARED
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+#else
+ if (devptr[0] != NULL)
+ __builtin_abort ();
+ if (devptr[1] != NULL)
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } }
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+ }
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
new file mode 100644
index 00000000000..4a1d5da98aa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90
@@ -0,0 +1,39 @@
+! { dg-do run }
+
+! Test no_create clause with data construct when data is present/not present.
+
+program no_create
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myvar, myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ myvar = 77
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .neqv. shared_memory) stop 10
+ if (acc_is_present (myarr) .neqv. shared_memory) stop 11
+ !$acc end data
+
+ !$acc enter data copyin (myvar, myarr)
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .eqv. .false.) stop 20
+ if (acc_is_present (myarr) .eqv. .false.) stop 21
+ !$acc end data
+ !$acc exit data copyout (myvar, myarr)
+
+ if (myvar .ne. 77) stop 30
+ do i = 1, n
+ if (myarr(i) .ne. 0) stop 31
+ end do
+end program no_create
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
new file mode 100644
index 00000000000..0b11f454aca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90
@@ -0,0 +1,90 @@
+! { dg-do run }
+
+! Test no_create clause with data/parallel constructs.
+
+program no_create
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myvar, myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ myvar = 55
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ call do_on_target(myvar, n, myarr)
+
+ if (shared_memory) then
+ if (myvar .ne. 44) stop 10
+ else
+ if (myvar .ne. 33) stop 11
+ end if
+ do i = 1, n
+ if (shared_memory) then
+ if (myarr(i) .ne. i * 2) stop 20
+ else
+ if (myarr(i) .ne. i) stop 21
+ end if
+ end do
+
+ myvar = 55
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc enter data copyin(myvar, myarr)
+ call do_on_target(myvar, n, myarr)
+ !$acc exit data copyout(myvar, myarr)
+
+ if (myvar .ne. 44) stop 30
+ do i = 1, n
+ if (myarr(i) .ne. i * 2) stop 31
+ end do
+end program no_create
+
+subroutine do_on_target (var, n, arr)
+ use openacc
+ implicit none
+ integer :: var, n, arr(n)
+ integer :: i
+
+!$acc data no_create (var, arr)
+
+if (acc_is_present(var)) 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 !no_create (var)
+ var = 44
+ !$acc end parallel
+else
+ var = 33
+end if
+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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90
new file mode 100644
index 00000000000..4362688e579
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90
@@ -0,0 +1,39 @@
+! { dg-do run }
+
+program main
+ use iso_c_binding, only: c_sizeof
+ use openacc, only: acc_is_present
+ implicit none
+ integer i
+ integer, parameter :: n = 100
+ real*4 b(n), c(n)
+ real :: d(n), e(n)
+ common /BLOCK/ d, e
+
+ !$acc enter data create(b) create(d)
+
+ if (.not. acc_is_present(b, c_sizeof(b))) stop 1
+ if (.not. acc_is_present(d, c_sizeof(d))) stop 2
+#if !ACC_MEM_SHARED
+ if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 3
+ if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(d))) stop 4
+#endif
+
+ !$acc parallel loop no_create(b) no_create(c) no_create(/BLOCK/)
+ do i = 1, n
+ b(i) = i
+ d(i) = -i
+ end do
+ !$acc end parallel loop
+
+ if (.not. acc_is_present(b, c_sizeof(b))) stop 5
+ if (.not. acc_is_present(d, c_sizeof(d))) stop 6
+#if !ACC_MEM_SHARED
+ if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 7
+ if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(e))) stop 8
+#endif
+
+ !$acc exit data copyout(b) copyout(d)
+ if (any(abs(b - [(real(i), i = 1, n)]) > 10*epsilon(b))) stop 9
+ if (any(abs(d - [(real(-i), i = 1, n)]) > 10*epsilon(d))) stop 10
+end program main
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [Patch] Add OpenACC 2.6's no_create
2019-12-18 12:41 ` Tobias Burnus
@ 2019-12-18 21:47 ` Thomas Schwinge
0 siblings, 0 replies; 12+ messages in thread
From: Thomas Schwinge @ 2019-12-18 21:47 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches, fortran, Jakub Jelinek, Julian Brown
[-- Attachment #1: Type: text/plain, Size: 5628 bytes --]
Hi Tobias!
On 2019-12-18T13:36:29+0100, Tobias Burnus <tobias@codesourcery.com> wrote:
> libgomp/target.c's gomp_map_vars_internal: it now uses the normal code
> path in the upper loop, except that one directly bails out when the
> 'key' has not been found (skipping the adjacent MAP_POINTER as well).
> The 'case' in the second loop is only reached, if tgt[i]->key == NULL
> (i.e. if not present) and one can unconditionally skip here. — This
> seems to be cleaner and should avoid some confusions :-)
Oh, great! It seems that you managed to de-cypher what my brain (or was
it my gut feeling?) told me to write down in these TODO comments that I
had added. ;-)
I have not now reviewed the details, but from the structure, your changes
looks good, and if it work, all the better.
I note you're building up a "dangerous" ;-) level of understanding of OMP
internals! :-)
> GOMP_MAP_POINTER, following MAP_IF_PRESENT: I am not sure about this.
So, what does a 'GOMP_MAP_POINTER' following a non-present
'GOMP_MAP_IF_PRESENT' mean -- is this 'GOMP_MAP_POINTER' operation
actually a no-op then, given that in the non-present case we'll just use
the host pointer? But if it is a no-op, should we then just let the
mapping code execute these 'GOMP_MAP_POINTER' operation, instead of
adding special-case code to skip them?
Are there any interactions with the OpenACC 2.6 manual deep copy
implementation maybe?
> The testsuite digests both mapping and skipping the map pointer. It
> looks a tad cleaner to avoid mapping the pointer (if the var is not
> present) – saving also few bytes and cpu cycles. On the down side, it
> adds an order dependence assumption, namely assuming that the
> MAP_POINTER after 'no_create'/MAP_IF_PRESENT always belongs to
> no_create. – [This patch follows the original patch and skips the
> map_pointer.]
Per his OpenACC 2.6 manual deep copy work, Julian has indeed established
that a 'GOMP_MAP_POINTER' is "only expected after some other mapping";
see "case GOMP_MAP_POINTER" in
<http://mid.mail-archive.com/65540b92dff74db1f15af930f87f7096d03e7efe.1576648001.git.julian@codesourcery.com>,
for example.
See also <https://gcc.gnu.org/wiki/LibgompPointerMappingKinds>
"unfinished notes on pointer mapping kinds" that Julian created.
The question then is, is it (a) correct (also per the OpenACC 2.6 manual
deep copy requirements) to skip these 'GOMP_MAP_POINTER' after
'GOMP_MAP_IF_PRESENT', and (b) only 'GOMP_MAP_POINTER' or also other
"variants", and/or (c) not do that skipping?
(For avoidance of doubt: this is fine to resolve later, given that it may
depend on the pending OpenACC 2.6 manual deep copy, and doesn't seem to
cause any issues at present.)
> Otherwise, except for added acc_is_present calls to no_create-3.c to
> check that no_create does not cause mapping and applying your/Thomas's
> patches, it matches my previous version, which was OK'ed. — Hence, I
> intent to commit it tomorrow, unless there are further comments.
ACK.
> On 12/17/19 8:11 PM, Tobias Burnus wrote:
>> On 12/3/19 4:16 PM, Thomas Schwinge wrote:
>>> Another thing: I've added just another little bit of testsuite
>>> coverage, and another thing broke. See "TODO" in attached incremental
>>> patch. […]
>> Files included, the other issue was XFAILed by you (and hence passed).
>> A fix for that issue is:
>> https://gcc.gnu.org/ml/gcc-patches/2019-12/msg01135.html — and a
>> completely separate issue. (That patch is small, very localized and
>> orthogonal to this patch.)
ACK, that's for later.
>>> The incremental Fortran test case changes have bene done in a rush; not
>>> sure if they make much sense, or should see some further work applied to
>>> them.
>>
>> I think one can do more, but they are fine. I am not 100% sure how to
>> read the following:
>>
>> ! 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.
(We inherited that from somebody else. I too didn't quickly understand
that.)
>> !$acc parallel !no_create (var)
>>
>> First, why is 'no_create(var)' now commented? – For this code, it
>> should really work both ways and independent whether commented boils
>> down to 'copy' (currently) or 'present' (with my other patch, linked
>> above).
If I remember correctly (remember: "done in a rush"), I think that was my
rationale: we should get kind-of an implicit 'no_create' here.
..., and then, learned something new this evening:
> .../testsuite/libgomp.oacc-fortran/no_create-1.f90 | 39 ++++++++++
> .../testsuite/libgomp.oacc-fortran/no_create-2.f90 | 90 ++++++++++++++++++++++
> .../testsuite/libgomp.oacc-fortran/no_create-3.F90 | 39 ++++++++++
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90
Why is this upper-case '.F90' when others are lower-case '.f90'?
> @@ -0,0 +1,39 @@
> +! { dg-do run }
> +
> +program main
> + use iso_c_binding, only: c_sizeof
> + use openacc, only: acc_is_present
> + implicit none
> + integer i
> + integer, parameter :: n = 100
> + real*4 b(n), c(n)
> + real :: d(n), e(n)
> + common /BLOCK/ d, e
> +
> + !$acc enter data create(b) create(d)
> +
> + if (.not. acc_is_present(b, c_sizeof(b))) stop 1
> + if (.not. acc_is_present(d, c_sizeof(d))) stop 2
> +#if !ACC_MEM_SHARED
> +[...]
Aha!
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 12+ messages in thread
end of thread, other threads:[~2019-12-18 21:06 UTC | newest]
Thread overview: 12+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <20181218132900.5ebf2f0f@squid.athome>
2018-12-19 21:31 ` [PATCH, og8] Add OpenACC 2.6 `no_create' clause support Maciej W. Rozycki
2018-12-20 14:36 ` Maciej W. Rozycki
2019-10-24 13:26 ` [Patch] Add OpenACC 2.6's no_create Tobias Burnus
2019-11-05 23:49 ` Thomas Schwinge
2019-11-06 12:43 ` Thomas Schwinge
2019-11-15 19:12 ` Tobias Burnus
2019-12-03 15:16 ` Thomas Schwinge
2019-12-03 17:39 ` Tobias Burnus
2019-12-13 12:07 ` Tobias Burnus
2019-12-17 19:25 ` Tobias Burnus
2019-12-18 12:41 ` Tobias Burnus
2019-12-18 21:47 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).