On 24/07/2020 8:27 am, Thomas Schwinge wrote:
That however completely defeats what we're intending to test here, which
is to "Test invalid intra-routine parallelism". The same problem has
been introduced in og10 commit 6a0b5806b24bfdefe0b0f3ccbcc51299e5195dca
"Various OpenACC reduction enhancements - test cases" for
'gcc/testsuite/c-c++-common/goacc/routine-4.c', which throughout changed:
-#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by
containing routine" }
+#pragma acc loop seq reduction (+:red)
Please revert that, and instead replace 'reduction (+:red)' with a
different "dummy loop operation" (just an empty loop body?), and in the
commit log state that this should've been included in the respective og10
commit adding the "gang reduction on an orphan loop" checking.
I have reverted all the previous changes and replaced the orphan loop gang
reductions with empty loops as suggested, and checked that the tests now pass.
Is this version okay for OG10?
Thanks
Kwok
From 280957dc80090bd0b92ad7a73f528851aad94051 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <k...@codesourcery.com>
Date: Sun, 26 Jul 2020 05:11:30 -0700
Subject: [PATCH] Fix c-c++-common/goacc/routine-4.c and
c-c++-common/goacc/routine-4-extern.c testcases
'Various OpenACC reduction enhancements - FE changes' (commit
6b3e1f7f05cd360bbd356b3f78511aa2ec3f40c3) introduced checks for gang
reductions on orphan loops. The checks triggered in the routine-4.c
and routine-4-extern.c testcases, requiring changes that effectively
rendered them useless as test cases.
This patch restores the original intent of the test cases, by restoring
the original tests and removing the orphan loop reductions that were
triggering the new check.
This patch should probably have been part of 'Various OpenACC reduction
enhancements - test cases' (commit 6a0b5806b24bfdefe0b0f3ccbcc51299e5195dca).
2020-07-26 Kwok Cheung Yeung <k...@codesourcery.com>
gcc/testsuite/
* c-c++-common/goacc/routine-4.c (seq, vector, worker, gang): Revert
previous changes. Remove loop reductions.
* c-c++-common/goacc/routine-4-extern.c (seq, vector, worker, gang):
Likewise.
---
.../c-c++-common/goacc/routine-4-extern.c | 72 ++++++++++------------
gcc/testsuite/c-c++-common/goacc/routine-4.c | 72 ++++++++++------------
2 files changed, 64 insertions(+), 80 deletions(-)
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-4-extern.c
b/gcc/testsuite/c-c++-common/goacc/routine-4-extern.c
index c23ddcf..ec44758 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-4-extern.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-4-extern.c
@@ -26,23 +26,21 @@ void seq (void)
extern_vector (); /* { dg-error "routine call uses" } */
extern_seq ();
- int red;
-
-#pragma acc loop reduction (+:red) // { dg-warning "insufficient partitioning"
}
+#pragma acc loop // { dg-warning "insufficient partitioning" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop gang reduction (+:red) // { dg-error "gang reduction on an
orphan loop" }
+#pragma acc loop gang // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by
containing routine" }
+#pragma acc loop worker // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red) // { dg-error "disallowed by
containing routine" }
+#pragma acc loop vector // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
void vector (void)
@@ -52,23 +50,21 @@ void vector (void)
extern_vector ();
extern_seq ();
- int red;
-
-#pragma acc loop reduction (+:red)
+#pragma acc loop
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop gang reduction (+:red) // { dg-error "gang reduction on an
orphan loop" }
+#pragma acc loop gang // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by
containing routine" }
+#pragma acc loop worker // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
void worker (void)
@@ -78,23 +74,21 @@ void worker (void)
extern_vector ();
extern_seq ();
- int red;
-
-#pragma acc loop reduction (+:red)
+#pragma acc loop
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop gang reduction (+:red) // { dg-error "gang reduction on an
orphan loop" }
+#pragma acc loop gang // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red)
+#pragma acc loop worker
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
void gang (void)
@@ -104,21 +98,19 @@ void gang (void)
extern_vector ();
extern_seq ();
- int red;
-
-#pragma acc loop reduction (+:red)
+#pragma acc loop
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop gang reduction (+:red) /* { dg-error "gang reduction on an
orphan loop" } */
+#pragma acc loop gang
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red)
+#pragma acc loop worker
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-4.c
b/gcc/testsuite/c-c++-common/goacc/routine-4.c
index ad17371..870ff64 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-4.c
@@ -17,23 +17,21 @@ void seq (void)
vector (); /* { dg-error "routine call uses" } */
seq ();
- int red;
-
-#pragma acc loop reduction (+:red) // { dg-warning "insufficient partitioning"
}
+#pragma acc loop // { dg-warning "insufficient partitioning" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop seq reduction (+:red)
+#pragma acc loop gang // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by
containing routine" }
+#pragma acc loop worker // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red) // { dg-error "disallowed by
containing routine" }
+#pragma acc loop vector // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
void vector (void) /* { dg-message "declared here" "1" } */
@@ -43,23 +41,21 @@ void vector (void) /* { dg-message "declared here" "1" } */
vector ();
seq ();
- int red;
-
-#pragma acc loop reduction (+:red)
+#pragma acc loop
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop seq reduction (+:red)
+#pragma acc loop gang // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by
containing routine" }
+#pragma acc loop worker // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
void worker (void) /* { dg-message "declared here" "2" } */
@@ -69,23 +65,21 @@ void worker (void) /* { dg-message "declared here" "2" } */
vector ();
seq ();
- int red;
-
-#pragma acc loop reduction (+:red)
+#pragma acc loop
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop seq reduction (+:red)
+#pragma acc loop gang // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red)
+#pragma acc loop worker
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
void gang (void) /* { dg-message "declared here" "3" } */
@@ -95,21 +89,19 @@ void gang (void) /* { dg-message "declared here" "3" } */
vector ();
seq ();
- int red;
-
-#pragma acc loop reduction (+:red)
+#pragma acc loop
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop seq reduction (+:red)
+#pragma acc loop gang
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop worker reduction (+:red)
+#pragma acc loop worker
for (int i = 0; i < 10; i++)
- red ++;
+ ;
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector
for (int i = 0; i < 10; i++)
- red ++;
+ ;
}
--
2.8.1