diff mbox

[committed] Fix zero based array map clause handling in #pragma omp target {data,update} (PR middle-end/64734)

Message ID 20150123182628.GY1746@tucnak.redhat.com
State New
Headers show

Commit Message

Jakub Jelinek Jan. 23, 2015, 6:26 p.m. UTC
Hi!

On the following testcase we ICE, because scan_sharing_clauses doesn't
install_var_field.  Generally in target {data,update} we don't want
to pointer translate when nothing will really use it, but for
OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION we actually do want data transfer
to happen and we have skipped doing it on the previous clause because
we rely on the next one to handle it.

Bootstrapped/regtested on x86_64-linux and i686-linux, committed.

2015-01-23  Jakub Jelinek  <jakub@redhat.com>

	PR middle-end/64734
	* omp-low.c (scan_sharing_clauses): Don't ignore
	OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION GOMP_MAP_POINTER clauses
	on target data/update constructs.

	* libgomp.c/pr64734.c: New test.


	Jakub
diff mbox

Patch

--- gcc/omp-low.c.jj	2015-01-19 14:40:46.000000000 +0100
+++ gcc/omp-low.c	2015-01-23 13:36:02.414003593 +0100
@@ -1834,7 +1834,8 @@  scan_sharing_clauses (tree clauses, omp_
 	      /* Ignore GOMP_MAP_POINTER kind for arrays in regions that are
 		 not offloaded; there is nothing to map for those.  */
 	      if (!is_gimple_omp_offloaded (ctx->stmt)
-		  && !POINTER_TYPE_P (TREE_TYPE (decl)))
+		  && !POINTER_TYPE_P (TREE_TYPE (decl))
+		  && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
 		break;
 	    }
 	  if (DECL_P (decl))
--- libgomp/testsuite/libgomp.c/pr64734.c.jj	2015-01-23 13:51:34.317722455 +0100
+++ libgomp/testsuite/libgomp.c/pr64734.c	2015-01-23 14:33:03.292638067 +0100
@@ -0,0 +1,55 @@ 
+/* PR middle-end/64734 */
+
+#include <stdlib.h>
+
+void
+foo (int *x, int *y)
+{
+  #pragma omp target map (alloc:x[0]) map (alloc:y[0:8])
+  {
+    int i;
+    for (i = 0; i < 8; i++)
+      if (y[i] != 2 + i)
+	break;
+    if (i != 8 || *x != 1)
+      *x = 6;
+    else
+      {
+	*x = 8;
+	for (i = 0; i < 8; i++)
+	  y[i] = 9 + i;
+      }
+  }
+  #pragma omp target update from (y[0:8]) from (x[0])
+}
+
+void
+bar (void)
+{
+  int x = 1, y[32] = { 0 };
+  #pragma omp target data map (to:y[0:32]) map (to:x)
+    ;
+}
+
+int
+main ()
+{
+  int x = 1, y[8] = { 2, 3, 4, 5, 6, 7, 8, 9 }, i;
+  #pragma omp target data map (to:y[0:8]) map (to:x)
+    ;
+  #pragma omp target data map (to:y[0:8]) map (to:x)
+    {
+      #pragma omp target update from (y[0:8]) from (x)
+    }
+
+  #pragma omp target data map (to:y[0:8]) map (to:x)
+    foo (&x, &y[0]);
+
+  if (x != 8)
+    abort ();
+  for (i = 0; i < 8; i++)
+    if (y[i] != 9 + i)
+      abort ();
+
+  return 0;
+}