diff mbox

[5/5] OpenACC tile clause support, libgomp testsuite patches

Message ID dfb007af-c533-3306-8b8e-df158b4be21d@mentor.com
State New
Headers show

Commit Message

Chung-Lin Tang Nov. 10, 2016, 10:47 a.m. UTC
Some additional tests and adjustments to existing ones were made.

2016-XX-XX  Nathan Sidwell  <nathan@codesourcery.com>
	    Chung-Lin Tang  <cltang@codesourcery.com>

        libgomp/
        * testsuite/libgomp.oacc-c-c++-common/tile-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust and
        add additional case.
	* testsuite/libgomp.oacc-c-c++-common/vprop.c: XFAIL under
        "openacc_nvidia_accel_selected".
        * libgomp.oacc-fortran/nested-function-1.f90 (test2): Add num_workers(8)
        clause.

Comments

Jakub Jelinek Nov. 11, 2016, 10:35 a.m. UTC | #1
On Thu, Nov 10, 2016 at 06:47:07PM +0800, Chung-Lin Tang wrote:
> Some additional tests and adjustments to existing ones were made.

> 

> 2016-XX-XX  Nathan Sidwell  <nathan@codesourcery.com>

> 	    Chung-Lin Tang  <cltang@codesourcery.com>

> 

>         libgomp/

>         * testsuite/libgomp.oacc-c-c++-common/tile-1.c: New.

> 	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust and

>         add additional case.

> 	* testsuite/libgomp.oacc-c-c++-common/vprop.c: XFAIL under

>         "openacc_nvidia_accel_selected".

>         * libgomp.oacc-fortran/nested-function-1.f90 (test2): Add num_workers(8)

>         clause.


Again, please fix up the whitespace in ChangeLog.  Ok with that fixed
(once the other patches are in of course).

	Jakub
diff mbox

Patch

Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c	(revision 241809)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c	(working copy)
@@ -112,7 +112,7 @@  int vector_1 (int *ary, int size)
 	ary[ix] = place ();
   }
 
-  return check (ary, size, 0, 0, 1);
+  return check (ary, size, 0, 1, 1);
 }
 
 int vector_2 (int *ary, int size)
@@ -196,10 +196,24 @@  int gang_3 (int *ary, int size)
 	ary[ix + jx * 64] = place ();
   }
 
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_4 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop auto
+    for (int jx = 0; jx <  size; jx++)
+      ary[jx] = place ();
+  }
+
   return check (ary, size, 1, 0, 1);
 }
 
-#define N (32*32*32)
+#define N (32*32*32*2)
 int main ()
 {
   int ondev = 0;
@@ -227,6 +241,8 @@  int main ()
     return 1;
   if (gang_3 (ary,  N))
     return 1;
+  if (gang_4 (ary,  N))
+    return 1;
 
   return 0;
 }
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c	(revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c	(revision 0)
@@ -0,0 +1,281 @@ 
+/* This code uses nvptx inline assembly guarded with acc_on_device, which is
+   not optimized away at -O0, and then confuses the target assembler.
+   { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
+
+/* { dg-additional-options "-fopenacc-dim=32" } */
+
+#include <stdio.h>
+#include <openacc.h>
+
+static int check (const int *ary, int size, int gp, int wp, int vp)
+{
+  int exit = 0;
+  int ix;
+  int gangs[32], workers[32], vectors[32];
+
+  for (ix = 0; ix < 32; ix++)
+    gangs[ix] = workers[ix] = vectors[ix] = 0;
+  
+  for (ix = 0; ix < size; ix++)
+    {
+      vectors[ary[ix] & 0xff]++;
+      workers[(ary[ix] >> 8) & 0xff]++;
+      gangs[(ary[ix] >> 16) & 0xff]++;
+    }
+
+  for (ix = 0; ix < 32; ix++)
+    {
+      if (gp)
+	{
+	  int expect = gangs[0];
+	  if (gangs[ix] != expect)
+	    {
+	      exit = 1;
+	      printf ("gang %d not used %d times\n", ix, expect);
+	    }
+	}
+      else if (ix && gangs[ix])
+	{
+	  exit = 1;
+	  printf ("gang %d unexpectedly used\n", ix);
+	}
+
+      if (wp)
+	{
+	  int expect = workers[0];
+	  if (workers[ix] != expect)
+	    {
+	      exit = 1;
+	      printf ("worker %d not used %d times\n", ix, expect);
+	    }
+	}
+      else if (ix && workers[ix])
+	{
+	  exit = 1;
+	  printf ("worker %d unexpectedly used\n", ix);
+	}
+
+      if (vp)
+	{
+	  int expect = vectors[0];
+	  if (vectors[ix] != expect)
+	    {
+	      exit = 1;
+	      printf ("vector %d not used %d times\n", ix, expect);
+	    }
+	}
+      else if (ix && vectors[ix])
+	{
+	  exit = 1;
+	  printf ("vector %d unexpectedly used\n", ix);
+	}
+      
+    }
+  return exit;
+}
+
+#pragma acc routine seq
+static int __attribute__((noinline)) place ()
+{
+  int r = 0;
+
+  if (acc_on_device (acc_device_nvidia))
+    {
+      int g = 0, w = 0, v = 0;
+
+      __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+      __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+      __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+      r = (g << 16) | (w << 8) | v;
+    }
+  return r;
+}
+
+static void clear (int *ary, int size)
+{
+  int ix;
+
+  for (ix = 0; ix < size; ix++)
+    ary[ix] = -1;
+}
+
+int gang_vector_1 (int *ary, int size)
+{
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(128) gang vector
+    for (int jx = 0; jx < size; jx++)
+      ary[jx] = place ();
+  }
+
+  return check (ary, size, 1, 0, 1);
+}
+
+int gang_vector_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) gang vector
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 1, 0, 1);
+}
+
+int gang_vector_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) gang vector
+    for (int jx = 0; jx < size; jx += 256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 1, 0, 1);
+}
+
+int worker_vector_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) worker vector
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 0, 1, 1);
+}
+
+int worker_vector_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) worker vector
+    for (int jx = 0; jx < size; jx += 256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 0, 1, 1);
+}
+
+int gang_worker_vector_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(32, 32)
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_worker_vector_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(32, 32)
+    for (int jx = 0; jx < size; jx += 256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_worker_vector_star_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(*, *)
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_worker_vector_star_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(*, *)
+    for (int jx = 0; jx < size; jx +=256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+#define N (32*32*32*8)
+int main ()
+{
+  int ondev = 0;
+
+#pragma acc parallel copy(ondev)
+  {
+    ondev = acc_on_device (acc_device_not_host);
+  }
+  if (!ondev)
+    return 0;
+  
+  int ary[N];
+  if (gang_vector_1 (ary, N))
+    return 1;
+  if (gang_vector_2a (ary, N))
+    return 1;
+  if (worker_vector_2a (ary, N))
+    return 1;
+  if (gang_worker_vector_2a (ary, N))
+    return 1;
+  if (gang_worker_vector_star_2a (ary, N))
+    return 1;
+  if (gang_vector_2b (ary, N))
+    return 1;
+  if (worker_vector_2b (ary, N))
+    return 1;
+  if (gang_worker_vector_2b (ary, N))
+    return 1;
+  if (gang_worker_vector_star_2b (ary, N))
+    return 1;
+  return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c	(revision 241809)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c	(working copy)
@@ -1,3 +1,6 @@ 
+/* { dg-do run } */
+/* { dg-xfail-run-if "PR78266" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
 #include <assert.h>
 
 #define test(type)				\
Index: libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90	(revision 241809)
+++ libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90	(working copy)
@@ -33,7 +33,7 @@  firstdo: do i = 1, 3
   subroutine test2
     integer :: a(3,3,3), k, kk, kkk, l, ll, lll
     a = 0
-    !$acc parallel
+    !$acc parallel num_workers(8)
     ! Use "gang(static:1)" here and below to effectively turn gang-redundant
     ! execution mode into something like gang-single.
     !$acc loop gang(static:1) collapse(1)