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)
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 <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@codesourcery.com>
Tobias Burnus <tob...@codesourcery.com>
Thomas Schwinge <tho...@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 <tho...@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
