diff mbox series

[og8] Add OpenACC 2.6 `no_create' clause support

Message ID alpine.DEB.2.21.9999.1812192056090.99920@build7-trusty-cs.sje.mentorg.com
State Accepted
Headers show
Series [og8] Add OpenACC 2.6 `no_create' clause support | expand

Commit Message

Maciej W. Rozycki Dec. 19, 2018, 9:30 p.m. UTC 
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

Comments

Maciej W. Rozycki Dec. 20, 2018, 2:30 p.m. UTC | #1 
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
diff mbox series

Patch

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