| Message ID | 5c6424d7-b85a-e3f2-c57f-f7e13a8aee83@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | libgomp, fortran: Apply if clause to all sub-constructs in combined OpenMP constructs | expand |
On Wed, Jun 24, 2020 at 05:47:06PM +0100, Kwok Cheung Yeung wrote: > There appears to be a bug in the handling of the 'if' clause (without a > directive name modifier) for combined OpenMP constructs in the Fortran > front-end: > > static void > gfc_split_omp_clauses (gfc_code *code, > gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM]) > { > ... > if (code->ext.omp_clauses != NULL) > { > if (mask & GFC_OMP_MASK_TARGET) > { > /* And this is copied to all. */ > clausesa[GFC_OMP_SPLIT_PARALLEL].if_expr > = code->ext.omp_clauses->if_expr; > } Yeah, looks like a pasto. > Currently, if 'target' is in the combined contruct, then the 'if' is applied > to the 'parallel' construct, but there are combined constructs with 'target' > but not 'parallel' (e.g. target teams distribute), which result in the 'if' > not getting applied at all. This is also redundant, as the unmodified if is > always applied to the 'parallel' construct if there is one. > > The patch changes the behaviour to match what the common C/C++ FE does, > which is to apply the 'if' to every applicable sub-construct in the combined > construct. I have included a testcase to check that the if clauses have been > applied correctly by the time it gets to the ME. I have also found a case > that results in an ICE (using 'target parallel' with an 'if' clause) - I > have commented out this out for now and filed it as PR 95869. > > I have tested for regressions in the gfortran and libgomp testsuites. Okay > for master/OG10? Ok, thanks. Though, while you are at it, if you'd like to also add the related, but still missing support for if (simd: ...), or if (cancel: ...) that C/C++ already do support, it would be greatly appreciated. Jakub
Hi Kwok, the TODO is fixed by the attached patch; I would be happy if you could handle this patch, e.g. together with your patch – or as follow up. (Lightly tested only, i.e. it fixes the ICE but I did not do a full testsuite run. But I regard it as obvious.) Tobias On 6/24/20 6:47 PM, Kwok Cheung Yeung wrote: > + ! TODO: This currently fails with an internal compiler error > + ! (PR 95869) > + !subroutine test_target_parallel > + ! do j = 1, N > + ! !$omp target parallel if(j .lt. LIMIT) map(tofrom: a(1:N)) > + ! do i = 1, N > + ! a(i) = a(i) + 1 > + ! end do > + ! !$omp end target parallel > + ! end do > + !end subroutine At least with my build (w/o your patch but with other patches), I see in the original dump: D.4049 = j <= 59; #pragma omp target map( ...) { { logical(kind=4) D.4049; #pragma omp parallel private(i) if(D.4049) { Namely, the assignment is in a different scope than the declaration of the variable. At the moment, I do not see why this fails – at a glance, it l Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On 24/06/2020 6:29 pm, Tobias Burnus wrote: > Hi Kwok, > > the TODO is fixed by the attached patch; I would be happy if you could handle > this patch, > e.g. together with your patch – or as follow up. > > (Lightly tested only, i.e. it fixes the ICE but I did not > do a full testsuite run. But I regard it as obvious.) Hello I have committed your patch along with the testcase as 'obvious'. I have confirmed that it does not regress the gfortran and libgomp testsuites. Kwok commit f530bac8a11da9c85bdd8e58d589747f9825e38d Author: Kwok Cheung Yeung <kcy@codesourcery.com> Date: Thu Jun 25 04:40:53 2020 -0700 fortran: Fix ICE when 'if' clause used with 'target parallel' (PR95869) 2020-06-25 Tobias Burnus <tobias@codesourcery.com> Kwok Cheung Yeung <kcy@codesourery.com> gcc/fortran/ PR fortran/95869 * trans-openmp.c (gfc_trans_omp_target): Use correct scoping block. gcc/testsuite/ PR fortran/95869 * gfortran.dg/gomp/combined-if.f90 (test_target_parallel): Re-enable. * gfortran.dg/gomp/pr95869.f90: New. diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 67b7094..22f8f96 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -5353,7 +5353,7 @@ gfc_trans_omp_target (gfc_code *code) pushlevel (); gfc_start_block (&iblock); tree inner_clauses - = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_PARALLEL], + = gfc_trans_omp_clauses (&iblock, &clausesa[GFC_OMP_SPLIT_PARALLEL], code->loc); stmt = gfc_trans_omp_code (code->block->next, true); stmt = build2_loc (input_location, OMP_PARALLEL, void_type_node, stmt, diff --git a/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 b/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 index 383086c..bf4a9a8 100644 --- a/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 @@ -18,17 +18,15 @@ contains end do end subroutine - ! TODO: This currently fails with an internal compiler error - ! (PR 95869) - !subroutine test_target_parallel - ! do j = 1, N - ! !$omp target parallel if(j .lt. LIMIT) map(tofrom: a(1:N)) - ! do i = 1, N - ! a(i) = a(i) + 1 - ! end do - ! !$omp end target parallel - ! end do - !end subroutine + subroutine test_target_parallel + do j = 1, N + !$omp target parallel if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + !$omp end target parallel + end do + end subroutine subroutine test_target_parallel_loop do j = 1, N @@ -105,6 +103,6 @@ contains end module -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target.* if\\(" 8 "omplower" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target.* if\\(" 9 "omplower" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp simd.* if\\(" 7 "omplower" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp parallel.* if\\(" 5 "omplower" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp parallel.* if\\(" 6 "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/pr95869.f90 b/gcc/testsuite/gfortran.dg/gomp/pr95869.f90 new file mode 100644 index 0000000..daa8d21 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr95869.f90 @@ -0,0 +1,18 @@ +! PR fortran/95869 +! { dg-do compile } + +program pr95869 + implicit none + + integer, parameter :: N = 100 + integer, parameter :: LIMIT = 60 + integer :: i, j + integer, dimension(N) :: a = (/ (i, i = 1,N) /) + do j = 1, N + !$omp target parallel if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + !$omp end target parallel + end do +end program
On Thu, 25 Jun 2020 at 15:24, Kwok Cheung Yeung <kcy@codesourcery.com> wrote: > > On 24/06/2020 6:29 pm, Tobias Burnus wrote: > > Hi Kwok, > > > > the TODO is fixed by the attached patch; I would be happy if you could handle > > this patch, > > e.g. together with your patch – or as follow up. > > > > (Lightly tested only, i.e. it fixes the ICE but I did not > > do a full testsuite run. But I regard it as obvious.) > > Hello > > I have committed your patch along with the testcase as 'obvious'. I have > confirmed that it does not regress the gfortran and libgomp testsuites. > Hi, I've noticed a regression since your commit, on arm aarch64 and x86: for instance on arm-linux-gnueabi: PASS: gfortran.dg/gomp/combined-if.f90 -O (test for excess errors) PASS: gfortran.dg/gomp/combined-if.f90 -O scan-tree-dump-times omplower "(?n)#pragma omp target.* if\\(" 9 gfortran.dg/gomp/combined-if.f90 -O : pattern found 4 times FAIL: gfortran.dg/gomp/combined-if.f90 -O scan-tree-dump-times omplower "(?n)#pragma omp simd.* if\\(" 7 PASS: gfortran.dg/gomp/combined-if.f90 -O scan-tree-dump-times omplower "(?n)#pragma omp parallel.* if\\(" 6 Not sure why you didn't see it? Thanks, Christophe > Kwok >
On 26/06/2020 8:35 am, Christophe Lyon wrote: > Hi, > > I've noticed a regression since your commit, on arm aarch64 and x86: > for instance on arm-linux-gnueabi: > PASS: gfortran.dg/gomp/combined-if.f90 -O (test for excess errors) > PASS: gfortran.dg/gomp/combined-if.f90 -O scan-tree-dump-times > omplower "(?n)#pragma omp target.* if\\(" 9 > gfortran.dg/gomp/combined-if.f90 -O : pattern found 4 times > FAIL: gfortran.dg/gomp/combined-if.f90 -O scan-tree-dump-times > omplower "(?n)#pragma omp simd.* if\\(" 7 > PASS: gfortran.dg/gomp/combined-if.f90 -O scan-tree-dump-times > omplower "(?n)#pragma omp parallel.* if\\(" 6 > > Not sure why you didn't see it? > Hello I was working with a compiler built with offloading support for nvptx, and that emits extra '#pragma omp simd' statements with an extra _simt_: #pragma omp simd _simduid_(simduid.21) _looptemp_(D.4127) _looptemp_(D.4128) _simt_ linear(i.13:1) linear(i:1) if(D.4108) #pragma omp simd _simduid_(simduid.22) _looptemp_(D.4127) _looptemp_(D.4128) linear(i.13:1) linear(i:1) if(D.4108) #pragma omp simd _simduid_(simduid.64) _looptemp_(D.4437) _looptemp_(D.4438) _simt_ linear(i.58:1) linear(i:1) if(D.3977) #pragma omp simd _simduid_(simduid.65) _looptemp_(D.4437) _looptemp_(D.4438) linear(i.58:1) linear(i:1) if(D.3977) #pragma omp simd _simduid_(simduid.106) _simt_ linear(i.101:1) linear(i:1) if(D.4681) #pragma omp simd _simduid_(simduid.107) linear(i.101:1) linear(i:1) if(D.4681) #pragma omp simd _simduid_(simduid.170) _looptemp_(D.5081) _looptemp_(D.5082) linear(i.164:1) linear(i:1) if(D.4069) If offloading is disabled, or targeted for amdgcn, the lines sith _simt_ do not appear. This appears to be controlled by omp_max_simt_vf() in omp-general.c, and is currently only non-zero for nvptx. I think the easiest fix would be to expect different numbers of matches depending on whether nvptx offloading is enabled. This requires an extra function in gcc/testsuite/lib/target-supports.exp. Okay for master/OG10? Thanks Kwok commit 04bdcaa20827d814c323847630c59ee843c51408 Author: Kwok Cheung Yeung <kcy@codesourcery.com> Date: Fri Jun 26 10:35:36 2020 -0700 Fix failure in gfortran.dg/gomp/combined-if.f90 test Enabling nvptx offloading results in extra '#pragma omp simd' statements in the tree dump with an extra '_simt_'. 2020-06-26 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/testsuite/ * testsuite/gfortran.dg/gomp/combined-if.f90: Adjust expected number of matches depending on whether nvptx offloading is supported. * testsuite/lib/target-supports.exp (check_effective_target_offload_nvptx): New. diff --git a/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 b/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 index bf4a9a8..0bb6c28 100644 --- a/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 @@ -104,5 +104,6 @@ contains end module ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target.* if\\(" 9 "omplower" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp simd.* if\\(" 7 "omplower" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp simd.* if\\(" 4 "omplower" { target { ! offload_nvptx } } } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp simd.* if\\(" 7 "omplower" { target { offload_nvptx } } } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp parallel.* if\\(" 6 "omplower" } } diff --git a/gcc/testsuite/lib/target-supports.exp b/gcc/testsuite/lib/target-supports.exp index cf0cfa1..2279361 100644 --- a/gcc/testsuite/lib/target-supports.exp +++ b/gcc/testsuite/lib/target-supports.exp @@ -9820,6 +9820,14 @@ proc check_effective_target_vect_max_reduc { } { return 0 } +# Return 1 if the compiler has been configured with nvptx offloading. + +proc check_effective_target_offload_nvptx { } { + return [check_no_compiler_messages offload_nvptx assembly { + int main () {return 0;} + } "-foffload=nvptx-none" ] +} + # Return 1 if the compiler has been configured with hsa offloading. proc check_effective_target_offload_hsa { } { @@ -9828,7 +9836,7 @@ proc check_effective_target_offload_hsa { } { } "-foffload=hsa" ] } -# Return 1 if the compiler has been configured with hsa offloading. +# Return 1 if the compiler has been configured with gcn offloading. proc check_effective_target_offload_gcn { } { return [check_no_compiler_messages offload_gcn assembly {
On Fri, Jun 26, 2020 at 06:55:30PM +0100, Kwok Cheung Yeung wrote: > Okay for master/OG10? Ok for master, not sure if I have anything to say for OG10. With a nit. > commit 04bdcaa20827d814c323847630c59ee843c51408 > Author: Kwok Cheung Yeung <kcy@codesourcery.com> > Date: Fri Jun 26 10:35:36 2020 -0700 > > Fix failure in gfortran.dg/gomp/combined-if.f90 test > > Enabling nvptx offloading results in extra '#pragma omp simd' statements > in the tree dump with an extra '_simt_'. > > 2020-06-26 Kwok Cheung Yeung <kcy@codesourcery.com> > > gcc/testsuite/ > * testsuite/gfortran.dg/gomp/combined-if.f90: Adjust expected number > of matches depending on whether nvptx offloading is supported. > * testsuite/lib/target-supports.exp > (check_effective_target_offload_nvptx): New. The testsuite/ prefix shouldn't be there, guess you wouldn't be able to commit it that way. Jakub
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 7e2f625..67b7094 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -4748,7 +4748,7 @@ gfc_split_omp_clauses (gfc_code *code, clausesa[GFC_OMP_SPLIT_TARGET].if_exprs[OMP_IF_TARGET] = code->ext.omp_clauses->if_exprs[OMP_IF_TARGET]; /* And this is copied to all. */ - clausesa[GFC_OMP_SPLIT_PARALLEL].if_expr + clausesa[GFC_OMP_SPLIT_TARGET].if_expr = code->ext.omp_clauses->if_expr; } if (mask & GFC_OMP_MASK_TEAMS) @@ -4832,6 +4832,9 @@ gfc_split_omp_clauses (gfc_code *code, /* Duplicate collapse. */ clausesa[GFC_OMP_SPLIT_SIMD].collapse = code->ext.omp_clauses->collapse; + /* And this is copied to all. */ + clausesa[GFC_OMP_SPLIT_SIMD].if_expr + = code->ext.omp_clauses->if_expr; } if (mask & GFC_OMP_MASK_TASKLOOP) { diff --git a/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 b/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 new file mode 100644 index 0000000..383086c --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/combined-if.f90 @@ -0,0 +1,110 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-omplower" } + +module combined_if + implicit none + + integer, parameter :: N = 100 + integer, parameter :: LIMIT = 60 + integer :: i, j + integer, dimension(N) :: a = (/ (i, i = 1,N) /) +contains + subroutine test_parallel_loop_simd + do j = 1, N + !$omp parallel do simd if(j .lt. LIMIT) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + + ! TODO: This currently fails with an internal compiler error + ! (PR 95869) + !subroutine test_target_parallel + ! do j = 1, N + ! !$omp target parallel if(j .lt. LIMIT) map(tofrom: a(1:N)) + ! do i = 1, N + ! a(i) = a(i) + 1 + ! end do + ! !$omp end target parallel + ! end do + !end subroutine + + subroutine test_target_parallel_loop + do j = 1, N + !$omp target parallel do if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + + subroutine test_target_parallel_loop_simd + do j = 1, N + !$omp target parallel do simd if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + + subroutine test_target_simd + do j = 1, N + !$omp target simd if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + + subroutine test_target_teams + do j = 1, N + !$omp target teams if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + !$omp end target teams + end do + end subroutine + + subroutine test_target_teams_distribute + do j = 1, N + !$omp target teams distribute if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + + subroutine test_target_teams_distibute_simd + do j = 1, N + !$omp target teams distribute simd if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + + subroutine test_target_teams_distribute_parallel_loop + do j = 1, N + !$omp target teams distribute parallel do if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + + subroutine test_target_teams_distribute_parallel_loop_simd + do j = 1, N + !$omp target teams distribute parallel do simd if(j .lt. LIMIT) map(tofrom: a(1:N)) + do i = 1, N + a(i) = a(i) + 1 + end do + end do + end subroutine + +end module + +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target.* if\\(" 8 "omplower" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp simd.* if\\(" 7 "omplower" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp parallel.* if\\(" 5 "omplower" } }
Hello There appears to be a bug in the handling of the 'if' clause (without a directive name modifier) for combined OpenMP constructs in the Fortran front-end: static void gfc_split_omp_clauses (gfc_code *code, gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM]) { ... if (code->ext.omp_clauses != NULL) { if (mask & GFC_OMP_MASK_TARGET) { /* And this is copied to all. */ clausesa[GFC_OMP_SPLIT_PARALLEL].if_expr = code->ext.omp_clauses->if_expr; } Currently, if 'target' is in the combined contruct, then the 'if' is applied to the 'parallel' construct, but there are combined constructs with 'target' but not 'parallel' (e.g. target teams distribute), which result in the 'if' not getting applied at all. This is also redundant, as the unmodified if is always applied to the 'parallel' construct if there is one. The patch changes the behaviour to match what the common C/C++ FE does, which is to apply the 'if' to every applicable sub-construct in the combined construct. I have included a testcase to check that the if clauses have been applied correctly by the time it gets to the ME. I have also found a case that results in an ICE (using 'target parallel' with an 'if' clause) - I have commented out this out for now and filed it as PR 95869. I have tested for regressions in the gfortran and libgomp testsuites. Okay for master/OG10? Thanks Kwok commit 052993de7457af85d5749b2ab119ffcc65e341e5 Author: Kwok Cheung Yeung <kcy@codesourcery.com> Date: Thu Jun 18 12:40:16 2020 -0700 libgomp, fortran: Apply if clause to all sub-constructs in combined OpenMP constructs The unmodified 'if' clause should be applied to all the sub-constructs that accept an 'if' clause in a combined OpenMP construct, and not just to the 'parallel' sub-construct. 2020-06-24 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/fortran/ * trans-openmp.c (gfc_split_omp_clauses): Add if clause to target and simd sub-constructs. gcc/testsuite/ * gfortran.dg/gomp/combined-if.f90: New.