diff mbox

[GOMP4,COMMITTED] OpenACC deviceptr clause.

Message ID 1401976816-10577-1-git-send-email-thomas@codesourcery.com
State New
Headers show

Commit Message

Thomas Schwinge June 5, 2014, 2 p.m. UTC 
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>

	gcc/c/
	* c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses):
	Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
	gcc/
	* gimplify.c (gimplify_scan_omp_clauses)
	(gimplify_adjust_omp_clauses): Handle
	OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
	* omp-low.c (scan_sharing_clauses, lower_oacc_offload)
	(lower_omp_target): Likewise.
	* tree-core.h (enum omp_clause_map_kind)
	<OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment.
	gcc/testsuite/
	* c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC
	deviceptr clause is now supported.
	* c-c++-common/goacc/deviceptr-1.c: Extend.
	* c-c++-common/goacc/deviceptr-2.c: New file.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211278 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                                 |  8 +++
 gcc/c/ChangeLog.gomp                               |  5 ++
 gcc/c/c-typeck.c                                   |  5 +-
 gcc/gimplify.c                                     |  7 ++-
 gcc/omp-low.c                                      | 60 +++++++++++++++++++---
 gcc/testsuite/ChangeLog.gomp                       |  5 ++
 .../c-c++-common/goacc/data-clause-duplicate-1.c   |  4 +-
 gcc/testsuite/c-c++-common/goacc/deviceptr-1.c     | 22 +++++++-
 gcc/testsuite/c-c++-common/goacc/deviceptr-2.c     | 23 +++++++++
 gcc/tree-core.h                                    |  3 +-
 10 files changed, 127 insertions(+), 15 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
diff mbox

Patch

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 7371aa5..88f09b3 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,13 @@ 
 2014-06-05  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* gimplify.c (gimplify_scan_omp_clauses)
+	(gimplify_adjust_omp_clauses): Handle
+	OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
+	* omp-low.c (scan_sharing_clauses, lower_oacc_offload)
+	(lower_omp_target): Likewise.
+	* tree-core.h (enum omp_clause_map_kind)
+	<OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment.
+
 	* gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>:
 	Don't block OMP_CLAUSE_MAP_FORCE_PRESENT.
 
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 91978db..1e80031 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,8 @@ 
+2014-06-05  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses):
+	Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
+
 2014-03-20  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* c-parser.c: Update comments.
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index c4ba531..839cdf7 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11747,6 +11747,7 @@  handle_omp_array_sections (tree c)
       OMP_CLAUSE_SIZE (c) = size;
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 	return false;
+      gcc_assert (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
       tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
       OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
       if (!c_mark_addressable (t))
@@ -12168,7 +12169,9 @@  c_finish_omp_clauses (tree clauses)
 	  else if (!c_mark_addressable (t))
 	    remove = true;
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-		     && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+		     && (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+			 || (OMP_CLAUSE_MAP_KIND (c)
+			     == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)))
 		   && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
diff --git gcc/gimplify.c gcc/gimplify.c
index 6eaf6fd..a1b6be6 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6015,7 +6015,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  switch (OMP_CLAUSE_MAP_KIND (c))
 	    {
 	    case OMP_CLAUSE_MAP_FORCE_DEALLOC:
-	    case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
 	      input_location = OMP_CLAUSE_LOCATION (c);
 	      /* TODO.  */
 	      sorry ("data clause not yet implemented");
@@ -6533,6 +6532,12 @@  gimplify_adjust_omp_clauses (tree *list_p)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
 		   && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER)
 	    {
+	      /* For OMP_CLAUSE_MAP_FORCE_DEVICEPTR, we'll never enter here,
+		 because for these, TREE_CODE (DECL_SIZE (decl)) will always be
+		 INTEGER_CST.  */
+	      gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+			  != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
 	      tree decl2 = DECL_VALUE_EXPR (decl);
 	      gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
 	      decl2 = TREE_OPERAND (decl2, 0);
diff --git gcc/omp-low.c gcc/omp-low.c
index 3e282c0..39f0598 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1708,6 +1708,18 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 		  && !POINTER_TYPE_P (TREE_TYPE (decl)))
 		break;
 	    }
+#if 0
+	  /* In target regions that are not offloaded, libgomp won't pay
+	     attention to OMP_CLAUSE_MAP_FORCE_DEVICEPTR -- but I think we need
+	     to handle it here anyway, in order to create a visible copy of the
+	     variable.  */
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+	    {
+	      if (!is_gimple_omp_offloaded (ctx->stmt))
+		break;
+	    }
+#endif
 	  if (DECL_P (decl))
 	    {
 	      if (DECL_SIZE (decl)
@@ -1723,6 +1735,10 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 		}
 	      else
 		{
+		  gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+			      || (OMP_CLAUSE_MAP_KIND (c)
+				  != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+			      || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
 		  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		      && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
 		      && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -1738,6 +1754,10 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	    {
 	      tree base = get_base_address (decl);
 	      tree nc = OMP_CLAUSE_CHAIN (c);
+	      gcc_assert (nc == NULL_TREE
+			  || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP
+			  || (OMP_CLAUSE_MAP_KIND (nc)
+			      != OMP_CLAUSE_MAP_FORCE_DEVICEPTR));
 	      if (DECL_P (base)
 		  && nc != NULL_TREE
 		  && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
@@ -1867,6 +1887,9 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	    }
 	  if (DECL_P (decl))
 	    {
+	      gcc_assert ((OMP_CLAUSE_MAP_KIND (c)
+			   != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+			  || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
 	      if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
 		  && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
 		  && !COMPLETE_TYPE_P (TREE_TYPE (decl)))
@@ -1878,6 +1901,9 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      else if (DECL_SIZE (decl)
 		       && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
 		{
+		  gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+			      != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
 		  tree decl2 = DECL_VALUE_EXPR (decl);
 		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
 		  decl2 = TREE_OPERAND (decl2, 0);
@@ -9100,6 +9126,10 @@  lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  {
 	    x = build_receiver_ref (var, true, ctx);
 	    tree new_var = lookup_decl (var, ctx);
+	    gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+			|| (OMP_CLAUSE_MAP_KIND (c)
+			    != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+			|| TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE);
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
 		&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -9199,6 +9229,10 @@  lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	      {
 		tree var = lookup_decl_in_outer_ctx (ovar, ctx);
 		tree x = build_sender_ref (ovar, ctx);
+		gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+			    || (OMP_CLAUSE_MAP_KIND (c)
+				!= OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+			    || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		    && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
 		    && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -9219,12 +9253,14 @@  lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		      = OMP_CLAUSE_MAP_KIND (c);
 		    if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
 			 && (map_kind & OMP_CLAUSE_MAP_TO))
-			|| map_kind == OMP_CLAUSE_MAP_POINTER)
+			|| map_kind == OMP_CLAUSE_MAP_POINTER
+			|| map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
 		      gimplify_assign (avar, var, &ilist);
 		    avar = build_fold_addr_expr (avar);
 		    gimplify_assign (x, avar, &ilist);
-		    if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
-			 && (map_kind & OMP_CLAUSE_MAP_FROM))
+		    if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+			  && (map_kind & OMP_CLAUSE_MAP_FROM))
+			 || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
 			&& !TYPE_READONLY (TREE_TYPE (var)))
 		      {
 			x = build_sender_ref (ovar, ctx);
@@ -10606,6 +10642,10 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  {
 	    x = build_receiver_ref (var, true, ctx);
 	    tree new_var = lookup_decl (var, ctx);
+	    gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+			|| (OMP_CLAUSE_MAP_KIND (c)
+			    != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+			|| TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE);
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
 		&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -10732,12 +10772,15 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	      {
 		tree var = lookup_decl_in_outer_ctx (ovar, ctx);
 		tree x = build_sender_ref (ovar, ctx);
+		gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+			    || (OMP_CLAUSE_MAP_KIND (c)
+				!= OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+			    || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		    && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
 		    && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
 		    && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
 		  {
-		    gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
 		    tree avar
 		      = create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL);
 		    mark_addressable (avar);
@@ -10747,19 +10790,20 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		  }
 		else if (is_gimple_reg (var))
 		  {
-		    gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
 		    tree avar = create_tmp_var (TREE_TYPE (var), NULL);
 		    mark_addressable (avar);
 		    enum omp_clause_map_kind map_kind
 		      = OMP_CLAUSE_MAP_KIND (c);
 		    if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
 			 && (map_kind & OMP_CLAUSE_MAP_TO))
-			|| map_kind == OMP_CLAUSE_MAP_POINTER)
+			|| map_kind == OMP_CLAUSE_MAP_POINTER
+			|| map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
 		      gimplify_assign (avar, var, &ilist);
 		    avar = build_fold_addr_expr (avar);
 		    gimplify_assign (x, avar, &ilist);
-		    if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
-			 && (map_kind & OMP_CLAUSE_MAP_FROM))
+		    if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+			  && (map_kind & OMP_CLAUSE_MAP_FROM))
+			 || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
 			&& !TYPE_READONLY (TREE_TYPE (var)))
 		      {
 			x = build_sender_ref (ovar, ctx);
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 4e0ee28..08ec907 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,10 @@ 
 2014-06-05  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC
+	deviceptr clause is now supported.
+	* c-c++-common/goacc/deviceptr-1.c: Extend.
+	* c-c++-common/goacc/deviceptr-2.c: New file.
+
 	* c-c++-common/goacc/data-clause-duplicate-1.c: Extend.
 	* c-c++-common/goacc/present-1.c: New file.
 
diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
index 5c5ab02..7a1cf68 100644
--- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
+++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -6,9 +6,7 @@  fun (void)
   ;
 #pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */
   ;
-#pragma acc data create(fp[:10]) deviceptr(fp)
-  /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */
-  /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */
+#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
   ;
 #pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
   ;
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 1ac63bd..cf2d809 100644
--- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -61,4 +61,24 @@  fun3 (void)
   ;
 }
 
-/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */
+extern struct s s1;
+extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */
+
+void
+fun4 (void)
+{
+  struct s *s1_p = &s1;
+  struct s *s2_p = &s2;
+
+#pragma acc parallel deviceptr(s1) /* { dg-error "'s1' is not a pointer variable" } */
+  ;
+
+#pragma acc parallel deviceptr(s2)
+  ;
+
+#pragma acc parallel deviceptr(s1_p)
+  s1_p = 0;
+
+#pragma acc parallel deviceptr(s2_p)
+  s2_p = 0;
+}
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-2.c gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
new file mode 100644
index 0000000..ac162b4
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
@@ -0,0 +1,23 @@ 
+void
+fun1 (void)
+{
+  char *a = 0;
+
+#pragma acc data deviceptr(a)
+  ++a;
+
+#pragma acc data deviceptr(a)
+#pragma acc parallel
+  ++a;
+
+#pragma acc data deviceptr(a)
+#pragma acc parallel deviceptr(a)
+  ++a;
+
+#pragma acc data
+#pragma acc parallel deviceptr(a)
+  ++a;
+
+#pragma acc parallel deviceptr(a)
+  ++a;
+}
diff --git gcc/tree-core.h gcc/tree-core.h
index 8603553..8b70c5b 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -1225,7 +1225,8 @@  enum omp_clause_map_kind
   OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL,
   /* Deallocate a mapping, without copying from device.  */
   OMP_CLAUSE_MAP_FORCE_DEALLOC,
-  /* Is a device pointer.  */
+  /* Is a device pointer.  OMP_CLAUSE_SIZE for these is unused; is implicitly
+     POINTER_SIZE / BITS_PER_UNIT.  */
   OMP_CLAUSE_MAP_FORCE_DEVICEPTR,
 
   /* End marker.  */