Add OpenACC 2.6's no_create

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-19  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>

Co-Authored-By: Maciej W. Rozycki <macro@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>

From-SVN: r279551

34 files changed, 616 insertions, 19 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 0f1a0b6..825b077 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,11 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC.
+	* tree-pretty-print.c (dump_omp_clause): Likewise.
+
 2019-12-18  Eric Botcazou  <ebotcazou@adacore.com>
 
 	* ira.c (ira): Use simple LRA algorithm when not optimizing.
diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index 966607f..77d928a 100644
--- a/gcc/c-family/ChangeLog
+++ b/gcc/c-family/ChangeLog
@@ -1,3 +1,11 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-pragma.h (pragma_omp_clause): Add
+	PRAGMA_OACC_CLAUSE_NO_CREATE.
+
 2019-12-17  Martin Sebor  <msebor@redhat.com>
 
 	PR c++/61339
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index bfe681b..3754c5f 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/ChangeLog b/gcc/c/ChangeLog
index 5371e9c..f4a088a 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,17 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* 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.
+
 2019-12-09  David Malcolm  <dmalcolm@redhat.com>
 
 	* c-objc-common.c (range_label_for_type_mismatch::get_text):
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index bfe5699..9b80088 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 36aedc0..ce5e649 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/ChangeLog b/gcc/cp/ChangeLog
index 90a6d18..05ca5ec 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,15 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* 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.
+
 2019-12-18  Paolo Carlini  <paolo.carlini@oracle.com>
 
 	* typeck.c (cxx_sizeof_or_alignof_type): Add location_t parameter
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 1f7526e..ce2e4b5 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -33622,7 +33622,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;
@@ -33988,7 +33990,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,
@@ -34024,6 +34029,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;
@@ -36586,6 +36594,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";
@@ -40391,6 +40403,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
@@ -40712,6 +40725,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)		\
@@ -40726,8 +40740,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)		\
@@ -40745,6 +40760,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 1d664af..69010dc 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5292,6 +5292,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/ChangeLog b/gcc/fortran/ChangeLog
index b3a5b0d..e45ec22 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,16 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* 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.
+
 2019-12-18  Harald Anlauf  <anlauf@gmx.de>
 
 	PR fortran/70853
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index f4a2b99..3907d14 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 dc0521b..576003d 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 b6da7b9..7153491 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 d422c20..deed83b 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/ChangeLog b/gcc/testsuite/ChangeLog
index 340f909..55a550c 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,15 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* 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.
+
 2019-12-18  Paolo Carlini  <paolo.carlini@oracle.com>
 
 	* g++.dg/diagnostic/alignof2.C: New.
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
index 228637f..6df5aa6 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 5d49f61..30c87a9 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 b94214e..30930a0 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 f16d62c..454417d 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 a70f1e7..5583ffb 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 2697bb7..e33653b 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 1cf7a91..6036173 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/ChangeLog b/include/ChangeLog
index 6364ab4..faeb5c4 100644
--- a/include/ChangeLog
+++ b/include/ChangeLog
@@ -1,3 +1,10 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC.
+
 2019-11-16  Tim Ruehsen  <tim.ruehsen@gmx.de>
 
 	* demangle.h (struct demangle_component): Add member
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 9e356cd..79c5de3 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/ChangeLog b/libgomp/ChangeLog
index 779f276..e1f6bb4 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,18 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+	    Maciej W. Rozycki  <macro@codesourcery.com>
+	    Tobias Burnus  <tobias@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* 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.
+
 2019-12-18  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* oacc-mem.c (goacc_enter_data): Refactor, so that it can be
diff --git a/libgomp/target.c b/libgomp/target.c
index 67cd80a..d83b353d 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -707,6 +707,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)
@@ -893,6 +908,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 0000000..22e0c20
--- /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 0000000..fbd01a2
--- /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 0000000..18466b8
--- /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 0000000..963cb3a
--- /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 0000000..6f0ace5
--- /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 0000000..4a1d5da
--- /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 0000000..0b11f45
--- /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 0000000..4362688
--- /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