On 10/22/2015 08:00 AM, Jakub Jelinek wrote:
> On Thu, Oct 22, 2015 at 07:47:01AM -0700, Cesar Philippidis wrote:
>>> But it is unclear from the parsing what from these is allowed:
>>
>> int v, w;
>> ...
>> gang(26) // equivalent to gang(num:26)
>> gang(v) // gang(num:v)
>> vector(length: 16) // vector(length: 16)
>> vector(length: v) // vector(length: v)
>> vector(16) // vector(length: 16)
>> vector(v) // vector(length: v)
>> worker(num: 16) // worker(num: 16)
>> worker(num: v) // worker(num: 16)
>> worker(16) // worker(num: 16)
>> worker(v) // worker(num: 16)
>> gang(16, 24) // technically gang(num:16, num:24) is acceptable but it
>> // should be an error
>> gang(v, w) // likewise
>> gang(static: 16, num: 5) // gang(static: 16, num: 5)
>> gang(static: v, num: w) // gang(static: v, num: w)
>> gang(num: 5, static: 4) // gang(num: 5, static: 4)
>> gang(num: v, static: w) // gang(num: v, static: w)
>>
>> Also note that the static argument can accept '*'.
>>
>>> and if the length: or num: part is really optional, then
>>> int length, num;
>>> vector(length)
>>> worker(num)
>>> gang(num, static: 6)
>>> gang(static: 5, num)
>>> should be also accepted (or subset thereof?).
>>
>> Interesting question. The spec is unclear. It defines gang, worker and
>> vector as follows in section 2.7 in the OpenACC 2.0a spec:
>>
>> gang [( gang-arg-list )]
>> worker [( [num:] int-expr )]
>> vector [( [length:] int-expr )]
>>
>> where gang-arg is one of:
>>
>> [num:] int-expr
>> static: size-expr
>>
>> and gang-arg-list may have at most one num and one static argument,
>> and where size-expr is one of:
>>
>> *
>> int-expr
>>
>> So I've interpreted that as a requirement that length and num must be
>> followed by an int-expr, whatever that is.
>
> My reading of the above is that
> vector(length)
> is equivalent to
> vector(length: length)
> and
> worker(num)
> is equivalent to
> vector(num: num)
> etc. Basically, neither length nor num aren't reserved identifiers,
> so you can use them for variable names, and if
> vector(v) is equivalent to vector(length: v), then
> vector(length) should be equivalent to vector(length:length)
> or
> vector(length + 1) should be equivalent to vector(length: length+1)
> static is a keyword that can't start an integral expression, so I guess
> it is fine if you issue an expected : diagnostics after it.
>
> In any case, please add a testcase (both C and C++) which covers all these
> allowed variants (ideally one testcase) and rejected variants (another
> testcase with dg-error).
>
> This is still an easy case, as even the C FE has 2 tokens lookup.
> E.g. for OpenMP map clause where
> map (always, tofrom: x)
> means one thing and
> map (always, tofrom, y)
> another one (map (tofrom: always, tofrom, y))
> I had to do quite ugly things to get around this.
Here are the updated test cases. Besides for adding a new test to
exercise the loop shape parsing, I also removed that assembly file
included in the original patch that Ilya noticed.
Is this OK for trunk?
Cesar
2015-10-23 Nathan Sidwell <nat...@codesourcery.com>
libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-g-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.
* 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.
2015-10-23 Cesar Philippidis <ce...@codesourcery.com>
gcc/testsuite/
* c-c++-common/goacc/loop-shape.c: New.
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
new file mode 100644
index 0000000..3cb3006
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
@@ -0,0 +1,197 @@
+/* Exercise *_parser_oacc_shape_clause by checking various combinations
+ of gang, worker and vector clause arguments. */
+
+/* { dg-compile } */
+
+int main ()
+{
+ int i;
+ int v, w;
+ int length, num;
+
+ /* Valid uses. */
+
+ #pragma acc kernels
+ #pragma acc loop gang worker vector
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(26)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length: 16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length: v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num: 16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num: v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: 16, num: 5)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: v, num: w)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, static: 6)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: 5, num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, static:*)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static:*, 1)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, static:*)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 5, static: 4)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: v, static: w)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, static:num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length:length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ /* Invalid uses. */
+
+ #pragma acc kernels
+#pragma acc loop gang(16, 24) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(v, w) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, num:5) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(length:num) /* { dg-error "expected .num. or .static." } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(5, length:length) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(num:length) /* { dg-error "expected .length. before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(length:5) /* { dg-error "expected .num. before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(1, num:2) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ return 0;
+}
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
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
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
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
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
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
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;
+}
