===================================================================
@@ -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;
}
===================================================================
@@ -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;
+}
===================================================================
@@ -1,3 +1,6 @@
+/* { dg-do run } */
+/* { dg-xfail-run-if "PR78266" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
#include <assert.h>
#define test(type) \
===================================================================
@@ -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)
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.