Re: [gomp4 06/14] omp-low: copy omp_data_o to shared memory on NVPTX

[Jakub Jelinek <jakub@redhat.com>]

On Tue, Nov 10, 2015 at 11:39:36AM +0100, Jakub Jelinek wrote:
> On Tue, Nov 03, 2015 at 05:25:53PM +0300, Alexander Monakov wrote:
> > Here's an alternative patch that does not depend on exposure of shared-memory
> > address space, and does not try to use pass_late_lower_omp.  It's based on
> > Bernd's suggestion to transform
> 
> FYI, I've committed a new testcase to gomp-4_5-branch that covers various
> target data sharing/team sharing/privatization parallel
> sharing/privatization offloading cases.

And another testcase, this time using only OpenMP 4.0 features, and trying
to test the behavior of addressable vars in declare target functions where
it is not clear if they are executed in teams, distribute or parallel for
contexts.

Wanted to look what LLVM generates here (tried llvm trunk), but they are
unable to parse #pragma omp distribute or #pragma omp declare target,
so it is hard to guess anything.

Tested with XeonPhi offloading as well as host fallback, committed to trunk.

2015-11-26  Jakub Jelinek  <jakub@redhat.com>

	* testsuite/libgomp.c/target-35.c: New test.

--- libgomp/testsuite/libgomp.c/target-35.c	(revision 0)
+++ libgomp/testsuite/libgomp.c/target-35.c	(working copy)
@@ -0,0 +1,129 @@
+#include <omp.h>
+#include <stdlib.h>
+
+#pragma omp declare target
+__attribute__((noinline))
+void
+foo (int x, int y, int z, int *a, int *b)
+{
+  if (x == 0)
+    {
+      int i, j;
+      for (i = 0; i < 64; i++)
+	#pragma omp parallel for shared (a, b)
+	for (j = 0; j < 32; j++)
+	  foo (3, i, j, a, b);
+    }
+  else if (x == 1)
+    {
+      int i, j;
+      #pragma omp distribute dist_schedule (static, 1)
+      for (i = 0; i < 64; i++)
+	#pragma omp parallel for shared (a, b)
+	for (j = 0; j < 32; j++)
+	  foo (3, i, j, a, b);
+    }
+  else if (x == 2)
+    {
+      int j;
+      #pragma omp parallel for shared (a, b)
+      for (j = 0; j < 32; j++)
+	foo (3, y, j, a, b);
+    }
+  else
+    {
+      #pragma omp atomic
+      b[y] += z;
+      #pragma omp atomic
+      *a += 1;
+    }
+}
+
+__attribute__((noinline))
+int
+bar (int x, int y, int z)
+{
+  int a, b[64], i;
+  a = 8;
+  for (i = 0; i < 64; i++)
+    b[i] = i;
+  foo (x, y, z, &a, b);
+  if (x == 0)
+    {
+      if (a != 8 + 64 * 32)
+	return 1;
+      for (i = 0; i < 64; i++)
+	if (b[i] != i + 31 * 32 / 2)
+	  return 1;
+    }
+  else if (x == 1)
+    {
+      int c = omp_get_num_teams ();
+      int d = omp_get_team_num ();
+      int e = d;
+      int f = 0;
+      for (i = 0; i < 64; i++)
+	if (i == e)
+	  {
+	    if (b[i] != i + 31 * 32 / 2)
+	      return 1;
+	    f++;
+	    e = e + c;
+	  }
+	else if (b[i] != i)
+	  return 1;
+      if (a < 8 || a > 8 + f * 32)
+	return 1;
+    }
+  else if (x == 2)
+    {
+      if (a != 8 + 32)
+	return 1;
+      for (i = 0; i < 64; i++)
+	if (b[i] != i + (i == y ? 31 * 32 / 2 : 0))
+	  return 1;
+    }
+  else if (x == 3)
+    {
+      if (a != 8 + 1)
+	return 1;
+      for (i = 0; i < 64; i++)
+	if (b[i] != i + (i == y ? z : 0))
+	  return 1;
+    }
+  return 0;
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+  int i, j, err = 0;
+  #pragma omp target map(tofrom:err)
+  #pragma omp teams reduction(+:err)
+  err += bar (0, 0, 0);
+  if (err)
+    abort ();
+  #pragma omp target map(tofrom:err)
+  #pragma omp teams reduction(+:err)
+  err += bar (1, 0, 0);
+  if (err)
+    abort ();
+  #pragma omp target map(tofrom:err)
+  #pragma omp teams reduction(+:err)
+  #pragma omp distribute
+  for (i = 0; i < 64; i++)
+    err += bar (2, i, 0);
+  if (err)
+    abort ();
+  #pragma omp target map(tofrom:err)
+  #pragma omp teams reduction(+:err)
+  #pragma omp distribute
+  for (i = 0; i < 64; i++)
+  #pragma omp parallel for reduction(+:err)
+    for (j = 0; j < 32; j++)
+      err += bar (3, i, j);
+  if (err)
+    abort ();
+  return 0;
+}


	Jakub