loop-g-1.c: New.

	* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: New.
	* testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: New.

From-SVN: r229487

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 0f6fd91de1f7a356b9ce6190f00533fa1cd059ce..8f44af02606b968b2f8039f730787869ce7a7d94 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,12 @@
+2015-10-28  Nathan Sidwell  <nathan@codesourcery.com>
+
+       * testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: New.
+       * testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: New.
+       * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: New.
+       * testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: New.
+       * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: New.
+       * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: New.
+
 2015-10-27  Nathan Sidwell  <nathan@codesourcery.com>
 
        * plugin/plugin-nvptx.c (nvptx_exec): Remove check on compute

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
new file mode 100644 (file)
index 0000000..58545d0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+  int ary[N];
+  int ix;
+  int exit = 0;
+  int ondev = 0;
+
+  for (ix = 0; ix < N;ix++)
+    ary[ix] = -1;
+  
+#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev)
+  {
+#pragma acc loop gang
+    for (unsigned ix = 0; ix < N; ix++)
+      {
+       if (__builtin_acc_on_device (5))
+         {
+           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));
+           ary[ix] = (g << 16) | (w << 8) | v;
+           ondev = 1;
+         }
+       else
+         ary[ix] = ix;
+      }
+  }
+
+  for (ix = 0; ix < N; ix++)
+    {
+      int expected = ix;
+      if(ondev)
+       {
+         int g = ix / ((N + 31) / 32);
+         int w = 0;
+         int v = 0;
+
+         expected = (g << 16) | (w << 8) | v;
+       }
+      
+      if (ary[ix] != expected)
+       {
+         exit = 1;
+         printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+       }
+    }
+  
+  return exit;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
new file mode 100644 (file)
index 0000000..c01c6fa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+  int ary[N];
+  int ix;
+  int exit = 0;
+  int ondev = 0;
+
+  for (ix = 0; ix < N;ix++)
+    ary[ix] = -1;
+  
+#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev)
+  {
+#pragma acc loop gang (static:1)
+    for (unsigned ix = 0; ix < N; ix++)
+      {
+       if (__builtin_acc_on_device (5))
+         {
+           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));
+           ary[ix] = (g << 16) | (w << 8) | v;
+           ondev = 1;
+         }
+       else
+         ary[ix] = ix;
+      }
+  }
+
+  for (ix = 0; ix < N; ix++)
+    {
+      int expected = ix;
+      if(ondev)
+       {
+         int g = ix % 32;
+         int w = 0;
+         int v = 0;
+
+         expected = (g << 16) | (w << 8) | v;
+       }
+      
+      if (ary[ix] != expected)
+       {
+         exit = 1;
+         printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+       }
+    }
+  
+  return exit;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
new file mode 100644 (file)
index 0000000..f23e2f3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
@@ -0,0 +1,59 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+  int ary[N];
+  int ix;
+  int exit = 0;
+  int ondev = 0;
+
+  for (ix = 0; ix < N;ix++)
+    ary[ix] = -1;
+  
+#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev)
+  {
+#pragma acc loop gang worker vector
+    for (unsigned ix = 0; ix < N; ix++)
+      {
+       if (__builtin_acc_on_device (5))
+         {
+           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));
+           ary[ix] = (g << 16) | (w << 8) | v;
+           ondev = 1;
+         }
+       else
+         ary[ix] = ix;
+      }
+  }
+
+  for (ix = 0; ix < N; ix++)
+    {
+      int expected = ix;
+      if(ondev)
+       {
+         int chunk_size = (N + 32*32*32 - 1) / (32*32*32);
+         
+         int g = ix / (chunk_size * 32 * 32);
+         int w = ix / 32 % 32;
+         int v = ix % 32;
+
+         expected = (g << 16) | (w << 8) | v;
+       }
+      
+      if (ary[ix] != expected)
+       {
+         exit = 1;
+         printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+       }
+    }
+  
+  return exit;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
new file mode 100644 (file)
index 0000000..70c6292
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+  int ary[N];
+  int ix;
+  int exit = 0;
+  int ondev = 0;
+
+  for (ix = 0; ix < N;ix++)
+    ary[ix] = -1;
+  
+#pragma acc parallel vector_length(32) copy(ary) copy(ondev)
+  {
+#pragma acc loop vector
+    for (unsigned ix = 0; ix < N; ix++)
+      {
+       if (__builtin_acc_on_device (5))
+         {
+           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));
+           ary[ix] = (g << 16) | (w << 8) | v;
+           ondev = 1;
+         }
+       else
+         ary[ix] = ix;
+      }
+  }
+
+  for (ix = 0; ix < N; ix++)
+    {
+      int expected = ix;
+      if(ondev)
+       {
+         int g = 0;
+         int w = 0;
+         int v = ix % 32;
+
+         expected = (g << 16) | (w << 8) | v;
+       }
+      
+      if (ary[ix] != expected)
+       {
+         exit = 1;
+         printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+       }
+    }
+  
+  return exit;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
new file mode 100644 (file)
index 0000000..5473c2d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+  int ary[N];
+  int ix;
+  int exit = 0;
+  int ondev = 0;
+
+  for (ix = 0; ix < N;ix++)
+    ary[ix] = -1;
+  
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+  {
+#pragma acc loop worker
+    for (unsigned ix = 0; ix < N; ix++)
+      {
+       if (__builtin_acc_on_device (5))
+         {
+           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));
+           ary[ix] = (g << 16) | (w << 8) | v;
+           ondev = 1;
+         }
+       else
+         ary[ix] = ix;
+      }
+  }
+
+  for (ix = 0; ix < N; ix++)
+    {
+      int expected = ix;
+      if(ondev)
+       {
+         int g = 0;
+         int w = ix % 32;
+         int v = 0;
+
+         expected = (g << 16) | (w << 8) | v;
+       }
+      
+      if (ary[ix] != expected)
+       {
+         exit = 1;
+         printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+       }
+    }
+  
+  return exit;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
new file mode 100644 (file)
index 0000000..85e4476
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+  int ary[N];
+  int ix;
+  int exit = 0;
+  int ondev = 0;
+
+  for (ix = 0; ix < N;ix++)
+    ary[ix] = -1;
+  
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+  {
+#pragma acc loop worker vector
+    for (unsigned ix = 0; ix < N; ix++)
+      {
+       if (__builtin_acc_on_device (5))
+         {
+           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));
+           ary[ix] = (g << 16) | (w << 8) | v;
+           ondev = 1;
+         }
+       else
+         ary[ix] = ix;
+      }
+  }
+
+  for (ix = 0; ix < N; ix++)
+    {
+      int expected = ix;
+      if(ondev)
+       {
+         int g = 0;
+         int w = (ix / 32) % 32;
+         int v = ix % 32;
+
+         expected = (g << 16) | (w << 8) | v;
+       }
+      
+      if (ary[ix] != expected)
+       {
+         exit = 1;
+         printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+       }
+    }
+  
+  return exit;
+}