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 <[email protected]> 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 <[email protected]> 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
