OpenACC update if_present runtime support

	gcc/
	* gimplify.c (gimplify_omp_target_update): Relax OpenACC update data
	mappings to GOMP_MAP_{TO,FROM} when the user specifies if_present.

	gcc/testsuite/
	* c-c++-common/goacc/update-if_present-1.c: Update test case.

	libgomp/
	* oacc-parallel.c (GOACC_update): Handle GOMP_MAP_{TO,FROM} for the
	if_present data clauses.
	* testsuite/libgomp.oacc-c-c++-common/update-2.c: New test.
	* testsuite/libgomp.oacc-fortran/update-3.f90: New test.


git-svn-id: https://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@247798 138bc75d-0d04-0410-961f-82ee72b054a4

6 files changed, 63 insertions, 1 deletions

diff --git a/gcc/ChangeLog.gomp b/gcc/ChangeLog.gomp
index b914cb1e65c..4468daadc04 100644
--- a/gcc/ChangeLog.gomp
+++ b/gcc/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2017-05-09  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* gimplify.c (gimplify_omp_target_update): Relax OpenACC update data
+	mappings to GOMP_MAP_{TO,FROM} when the user specifies if_present.
+
 2017-05-05  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* tree-nested.c (convert_nonlocal_omp_clauses)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index af908f44496..47fe9eeb3eb 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10034,6 +10034,25 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 			     ort, TREE_CODE (expr));
   gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr),
 			       TREE_CODE (expr));
+  if (TREE_CODE (expr) == OACC_UPDATE
+      && find_omp_clause (OMP_STANDALONE_CLAUSES (expr), OMP_CLAUSE_IF_PRESENT))
+    {
+      /* The runtime uses GOMP_MAP_{TO,FROM} to denote the if_present
+	 clause.  */
+      for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	  switch (OMP_CLAUSE_MAP_KIND (c))
+	    {
+	    case GOMP_MAP_FORCE_TO:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+	      break;
+	    case GOMP_MAP_FORCE_FROM:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FROM);
+	      break;
+	    default:
+	      break;
+	    }
+    }
   stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
 
   gimplify_seq_add_stmt (pre_p, stmt);
diff --git a/gcc/testsuite/ChangeLog.gomp b/gcc/testsuite/ChangeLog.gomp
index d1f7c8430f6..952b1010482 100644
--- a/gcc/testsuite/ChangeLog.gomp
+++ b/gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,7 @@
+2017-05-09  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* c-c++-common/goacc/update-if_present-1.c: Update test case.
+
 2017-05-05  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* g++.dg/goacc/update-1.C: Update.
diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
index 519393cf2a9..5a19ee132c4 100644
--- a/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
@@ -12,6 +12,18 @@ t ()
 #pragma acc update device(b) async if_present
 #pragma acc update host(c[1:3]) wait(4) if_present
 #pragma acc update self(c) device(b) host (a) async(10) if (a == 5) if_present
+
+#pragma acc update self(a)
+#pragma acc update device(b) async
+#pragma acc update host(c[1:3]) wait(4)
+#pragma acc update self(c) device(b) host (a) async(10) if (a == 5)
 }
 
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_update if_present" 4 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present map.from:a .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present async.-1. map.to:b .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present wait.4. map.from:c.1. .len: 12.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present if.D...... async.10. map.from:a .len: 4.. map.to:b .len: 4.. map.from:c .len: 40.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update map.force_from:a .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update async.-1. map.force_to:b .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update wait.4. map.force_from:c.1. .len: 12.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if.D...... async.10. map.force_from:a .len: 4.. map.force_to:b .len: 4.. map.force_from:c .len: 40.." 1 "omplower" } } */
diff --git a/libgomp/ChangeLog.gomp b/libgomp/ChangeLog.gomp
index 5ea5c77e5c3..18c32afa3fe 100644
--- a/libgomp/ChangeLog.gomp
+++ b/libgomp/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2017-05-09  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* oacc-parallel.c (GOACC_update): Handle GOMP_MAP_{TO,FROM} for the
+	if_present data clauses.
+	* testsuite/libgomp.oacc-c-c++-common/update-2.c: New test.
+	* testsuite/libgomp.oacc-fortran/update-3.f90: New test.
+
 2017-05-03  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust test case
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 66acdf6bb3c..de70ac0b6aa 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -683,14 +683,29 @@ GOACC_update (int device, size_t mapnum,
 
 	      /* Restore the host pointer.  */
 	      *(uintptr_t *) hostaddrs[i] = t;
+	      update_device = false;
 	    }
 	  break;
 
+	case GOMP_MAP_TO:
+	  if (!acc_is_present (hostaddrs[i], sizes[i]))
+	    {
+	      update_device = false;
+	      break;
+	    }
+	  /* Fallthru  */
 	case GOMP_MAP_FORCE_TO:
 	  update_device = true;
 	  acc_update_device (hostaddrs[i], sizes[i]);
 	  break;
 
+	case GOMP_MAP_FROM:
+	  if (!acc_is_present (hostaddrs[i], sizes[i]))
+	    {
+	      update_device = false;
+	      break;
+	    }
+	  /* Fallthru  */
 	case GOMP_MAP_FORCE_FROM:
 	  update_device = false;
 	  acc_update_self (hostaddrs[i], sizes[i]);