This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4, committed] Fix OACC_LOOP usage in goacc tests
- From: Tom de Vries <Tom_deVries at mentor dot com>
- To: Chung-Lin Tang <cltang at codesourcery dot com>
- Cc: "gcc-patches at gnu dot org" <gcc-patches at gnu dot org>, Jakub Jelinek <jakub at redhat dot com>
- Date: Sat, 18 Jul 2015 09:05:11 +0200
- Subject: [gomp4, committed] Fix OACC_LOOP usage in goacc tests
- Authentication-results: sourceware.org; auth=none
Hi Chung-Lin,
when compiling f.i. kernels-loop-acc-loop.c with -Wall, I ran into:
...
In file included from kernels-loop-acc-loop.c:8:0:
kernels-loop.c: In function âmainâ:
kernels-loop.c:30:0: warning: ignoring #pragma ACC_LOOP [-Wunknown-pragmas]
...
kernels-loop-acc-loop.c contains:
...
#define ACC_LOOP acc loop
#include "kernels-loop.c"
...
and kernels-loop.c contains:
...
#pragma ACC_LOOP
...
Unfortunately, this expands to '#pragma ACC_LOOP', instead of the
desired '#pragma acc loop'.
Something that works, is:
...
#define ACC_LOOP loop
#pragma acc ACC_LOOP
...
But the empty variant (the one we need for kernels-loop.c) doesn't work.
We either get 'ignoring #pragma ACC_LOOP' for:
...
#define ACC_LOOP
#pragma acc ACC_LOOP
...
or 'ignoring #pragma acc' for:
...
#define ACC_LOOP ""
#pragma acc ACC_LOOP
...
It would be nice if the openacc standard defined a pragma acc nop.
For now, I've fixed this using simple conditional compilation:
...
#ifdef ACC_LOOP
#pragma acc loop
#endif
...
which works for both:
- undefined ACC_LOOP, and
- #define ACC_LOOP
[ Btw, note that this scheme:
...
#ifdef ACC_LOOP
#pragma acc ACC_LOOP
#endif
...
would work with:
- undefined ACC_LOOP,
- #define ACC_LOOP loop
- #define ACC_LOOP loop independent ]
Committed to gomp-4_0-branch.
Thanks,
- Tom
Fix OACC_LOOP usage in goacc tests
2015-07-17 Tom de Vries <tom@codesourcery.com>
* c-c++-common/goacc/kernels-loop-2-acc-loop.c (ACC_LOOP): Define empty.
* c-c++-common/goacc/kernels-loop-3-acc-loop.c: Same.
* c-c++-common/goacc/kernels-loop-acc-loop.c: Same.
* c-c++-common/goacc/kernels-loop-n-acc-loop.c: Same.
* c-c++-common/goacc/kernels-loop-2.c (ACC_LOOP): Remove default define.
(main): Define #pragma acc loop conditional on ACC_LOOP.
* c-c++-common/goacc/kernels-loop-3.c: Same.
* c-c++-common/goacc/kernels-loop.c: Same.
* c-c++-common/goacc/kernels-loop-n.c (ACC_LOOP): Remove default define.
(foo): Define #pragma acc loop conditional on ACC_LOOP.
---
.../c-c++-common/goacc/kernels-loop-2-acc-loop.c | 2 +-
gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c | 16 +++++++++-------
.../c-c++-common/goacc/kernels-loop-3-acc-loop.c | 2 +-
gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c | 8 +++-----
gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c | 2 +-
.../c-c++-common/goacc/kernels-loop-n-acc-loop.c | 2 +-
gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c | 8 +++-----
gcc/testsuite/c-c++-common/goacc/kernels-loop.c | 8 +++-----
8 files changed, 22 insertions(+), 26 deletions(-)
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-2-acc-loop.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-2-acc-loop.c
index abfcb85..c041ca5 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-2-acc-loop.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-2-acc-loop.c
@@ -4,7 +4,7 @@
/* { dg-additional-options "-fdump-tree-optimized" } */
/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
-#define ACC_LOOP acc loop
+#define ACC_LOOP
#include "kernels-loop-2.c"
/* Check that only three loops are analyzed, and that all can be
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c
index e175b5b..d3f8805 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c
@@ -8,10 +8,6 @@
#define N (1024 * 512)
#define COUNTERTYPE unsigned int
-#ifndef ACC_LOOP
-#define ACC_LOOP
-#endif
-
int
main (void)
{
@@ -25,21 +21,27 @@ main (void)
#pragma acc kernels copyout (a[0:N])
{
- #pragma ACC_LOOP
+#ifdef ACC_LOOP
+ #pragma acc loop
+#endif
for (COUNTERTYPE i = 0; i < N; i++)
a[i] = i * 2;
}
#pragma acc kernels copyout (b[0:N])
{
- #pragma ACC_LOOP
+#ifdef ACC_LOOP
+ #pragma acc loop
+#endif
for (COUNTERTYPE i = 0; i < N; i++)
b[i] = i * 4;
}
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
{
- #pragma ACC_LOOP
+#ifdef ACC_LOOP
+ #pragma acc loop
+#endif
for (COUNTERTYPE ii = 0; ii < N; ii++)
c[ii] = a[ii] + b[ii];
}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-3-acc-loop.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-3-acc-loop.c
index d6447a1..7e748c0 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-3-acc-loop.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-3-acc-loop.c
@@ -4,7 +4,7 @@
/* { dg-additional-options "-fdump-tree-optimized" } */
/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
-#define ACC_LOOP acc loop
+#define ACC_LOOP
#include "kernels-loop-3.c"
/* Check that only one loop is analyzed, and that it can be parallelized. */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c
index 7d8848c..2620075 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c
@@ -8,10 +8,6 @@
#define N (1024 * 512)
#define COUNTERTYPE unsigned int
-#ifndef ACC_LOOP
-#define ACC_LOOP
-#endif
-
int
main (void)
{
@@ -26,7 +22,9 @@ main (void)
#pragma acc kernels copy (c[0:N])
{
- #pragma ACC_LOOP
+#ifdef ACC_LOOP
+ #pragma acc loop
+#endif
for (COUNTERTYPE ii = 0; ii < N; ii++)
c[ii] = c[ii] + ii + 1;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c
index 0e44f7e..5b4e860 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c
@@ -4,7 +4,7 @@
/* { dg-additional-options "-fdump-tree-optimized" } */
/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
-#define ACC_LOOP acc loop
+#define ACC_LOOP
#include "kernels-loop.c"
/* Check that only one loop is analyzed, and that it can be parallelized. */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-n-acc-loop.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-n-acc-loop.c
index af8b3e9..5e4a993 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-n-acc-loop.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-n-acc-loop.c
@@ -4,7 +4,7 @@
/* { dg-additional-options "-fdump-tree-optimized" } */
/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
-#define ACC_LOOP acc loop
+#define ACC_LOOP
#include "kernels-loop-n.c"
/* Check that only one loop is analyzed, and that it can be parallelized. */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c
index 861f50f..411959e 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c
@@ -8,10 +8,6 @@
#define N ((1024 * 512) + 1)
#define COUNTERTYPE unsigned int
-#ifndef ACC_LOOP
-#define ACC_LOOP
-#endif
-
int
foo (COUNTERTYPE n)
{
@@ -31,7 +27,9 @@ foo (COUNTERTYPE n)
#pragma acc kernels copyin (a[0:n], b[0:n]) copyout (c[0:n])
{
- #pragma ACC_LOOP
+#ifdef ACC_LOOP
+ #pragma acc loop
+#endif
for (COUNTERTYPE ii = 0; ii < n; ii++)
c[ii] = a[ii] + b[ii];
}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop.c
index 877176e..ecc5aef 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-loop.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop.c
@@ -8,10 +8,6 @@
#define N (1024 * 512)
#define COUNTERTYPE unsigned int
-#ifndef ACC_LOOP
-#define ACC_LOOP
-#endif
-
int
main (void)
{
@@ -31,7 +27,9 @@ main (void)
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
{
- #pragma ACC_LOOP
+#ifdef ACC_LOOP
+ #pragma acc loop
+#endif
for (COUNTERTYPE ii = 0; ii < N; ii++)
c[ii] = a[ii] + b[ii];
}
--
1.9.1