* OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] [not found] <20190508145157.08beb4df@squid.athome> @ 2022-01-13 9:54 ` Thomas Schwinge 2022-03-01 16:46 ` Thomas Schwinge ` (2 more replies) 0 siblings, 3 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-01-13 9:54 UTC (permalink / raw) To: Julian Brown, gcc-patches, fortran; +Cc: asolokha [-- Attachment #1: Type: text/plain, Size: 1352 bytes --] Hi! On 2019-05-08T14:51:57+0100, Julian Brown <julian@codesourcery.com> wrote: > - The "addressable" bit is set during the kernels conversion pass for > variables that have "create" (alloc) clauses created for them in the > synthesised outer data region (instead of in the front-end, etc., > where it can't be done accurately). Such variables actually have > their address taken during transformations made in a later pass > (omp-low, I think), but there's a phase-ordering problem that means > the flag should be set earlier. The actual issue is a bit different, but yes, there is a problem. The related ICE has also been reported as <https://gcc.gnu.org/PR100280> "ICE in lower_omp_target, at omp-low.c:12287". (And I'm confused why we didn't run into that with the OpenACC 'kernels' decomposition originally.) I've pushed to master branch commit 9b32c1669aad5459dd053424f9967011348add83 "OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280]", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-OpenACC-kernels-decomposition-Mark-variables-used-in.patch --] [-- Type: text/x-diff, Size: 39847 bytes --] From 9b32c1669aad5459dd053424f9967011348add83 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Thu, 16 Dec 2021 22:02:37 +0100 Subject: [PATCH] OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] ... as otherwise 'gcc/omp-low.c:lower_omp_target' has to create a temporary: 13073 else if (is_gimple_reg (var)) 13074 { 13075 gcc_assert (offloaded); 13076 tree avar = create_tmp_var (TREE_TYPE (var)); 13077 mark_addressable (avar); ..., which (a) is only implemented for actualy *offloaded* regions (but not data regions), and (b) the subsequently synthesized code for writing to and later reading back from the temporary fundamentally conflicts with OpenACC 'async' (as used by OpenACC 'kernels' decomposition). That's all not trivial to make work, so let's just avoid this case. gcc/ PR middle-end/100280 * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region): Mark variables used in synthesized data clauses as addressable. gcc/testsuite/ PR middle-end/100280 * c-c++-common/goacc/kernels-decompose-pr100280-1.c: New. * c-c++-common/goacc/classify-kernels-parloops.c: Likewise. * c-c++-common/goacc/classify-kernels-unparallelized-parloops.c: Likewise. * c-c++-common/goacc/classify-kernels-unparallelized.c: Test '--param openacc-kernels=decompose'. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/kernels-decompose-2.c: Update. * c-c++-common/goacc/kernels-decompose-ice-1.c: Remove. * c-c++-common/goacc/kernels-decompose-ice-2.c: Likewise. * gfortran.dg/goacc/classify-kernels-parloops.f95: New. * gfortran.dg/goacc/classify-kernels-unparallelized-parloops.f95: Likewise. * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Test '--param openacc-kernels=decompose'. * gfortran.dg/goacc/classify-kernels.f95: Likewise. libgomp/ PR middle-end/100280 * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. Suggested-by: Julian Brown <julian@codesourcery.com> --- gcc/omp-oacc-kernels-decompose.cc | 6 +- .../goacc/classify-kernels-parloops.c | 41 +++++++ ...classify-kernels-unparallelized-parloops.c | 45 +++++++ .../goacc/classify-kernels-unparallelized.c | 5 +- .../c-c++-common/goacc/classify-kernels.c | 5 +- .../c-c++-common/goacc/kernels-decompose-2.c | 16 ++- .../goacc/kernels-decompose-ice-1.c | 114 ------------------ .../goacc/kernels-decompose-ice-2.c | 22 ---- .../goacc/kernels-decompose-pr100280-1.c | 19 +++ .../goacc/classify-kernels-parloops.f95 | 43 +++++++ ...assify-kernels-unparallelized-parloops.f95 | 47 ++++++++ .../goacc/classify-kernels-unparallelized.f95 | 5 +- .../gfortran.dg/goacc/classify-kernels.f95 | 5 +- .../declare-vla-kernels-decompose-ice-1.c | 2 +- .../libgomp.oacc-c-c++-common/f-asyncwait-1.c | 53 ++++---- .../kernels-decompose-1.c | 6 +- 16 files changed, 264 insertions(+), 170 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-kernels-parloops.c create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized-parloops.c delete mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c delete mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-kernels-parloops.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized-parloops.f95 diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index cf9718f55a5..4ca899d5ece 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -793,7 +793,8 @@ make_data_region_try_statement (location_t loc, gimple *body) /* If INNER_BIND_VARS holds variables, build an OpenACC data region with location LOC containing BODY and having 'create (var)' clauses for each - variable. If INNER_CLEANUP is present, add a try-finally statement with + variable (as a side effect, such variables also get TREE_ADDRESSABLE set). + If INNER_CLEANUP is present, add a try-finally statement with this cleanup code in the finally block. Return the new data region, or the original BODY if no data region was needed. */ @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, inner_data_clauses = new_clause; prev_mapped_var = v; + + /* See <https://gcc.gnu.org/PR100280>. */ + TREE_ADDRESSABLE (v) = 1; } } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-parloops.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-parloops.c new file mode 100644 index 00000000000..f3685f2e8c5 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-parloops.c @@ -0,0 +1,41 @@ +/* Check offloaded function's attributes and classification for OpenACC + kernels. */ + +/* { dg-additional-options "--param openacc-kernels=parloops" } */ + +/* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } + { dg-additional-options "-fdump-tree-ompexp" } + { dg-additional-options "-fdump-tree-parloops1-all" } + { dg-additional-options "-fdump-tree-oaccloops" } */ + +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + +#define N 1024 + +extern unsigned int *__restrict a; +extern unsigned int *__restrict b; +extern unsigned int *__restrict c; + +void KERNELS () +{ +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + for (unsigned int i = 0; i < N; i++) + c[i] = a[i] + b[i]; +} + +/* Check the offloaded function's attributes. + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels, omp target entrypoint\\)\\)" 1 "ompexp" } } */ + +/* Check that exactly one OpenACC kernels construct is analyzed, and that it + can be parallelized. + { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops1" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "parloops1" } } + { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } */ + +/* Check the offloaded function's classification and compute dimensions (will + always be 1 x 1 x 1 for non-offloading compilation). + { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized-parloops.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized-parloops.c new file mode 100644 index 00000000000..6522caf9135 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized-parloops.c @@ -0,0 +1,45 @@ +/* Check offloaded function's attributes and classification for unparallelized + OpenACC kernels. */ + +/* { dg-additional-options "--param openacc-kernels=parloops" } */ + +/* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } + { dg-additional-options "-fdump-tree-ompexp" } + { dg-additional-options "-fdump-tree-parloops1-all" } + { dg-additional-options "-fdump-tree-oaccloops" } */ + +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + +#define N 1024 + +extern unsigned int *__restrict a; +extern unsigned int *__restrict b; +extern unsigned int *__restrict c; + +/* An "extern"al mapping of loop iterations/array indices makes the loop + unparallelizable. */ +extern unsigned int f (unsigned int); + +void KERNELS () +{ +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (unsigned int i = 0; i < N; i++) + c[i] = a[f (i)] + b[f (i)]; +} + +/* Check the offloaded function's attributes. + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels, omp target entrypoint\\)\\)" 1 "ompexp" } } */ + +/* Check that exactly one OpenACC kernels construct is analyzed, and that it + can't be parallelized. + { dg-final { scan-tree-dump-times "FAILED:" 1 "parloops1" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "parloops1" } } + { dg-final { scan-tree-dump-not "SUCCESS: may be parallelized" "parloops1" } } */ + +/* Check the offloaded function's classification and compute dimensions (will + always be 1 x 1 x 1 for non-offloading compilation). + { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index 1d12658790d..daa8fcb7662 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -1,8 +1,10 @@ /* Check offloaded function's attributes and classification for unparallelized OpenACC kernels. */ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + /* { dg-additional-options "-O2" } - { dg-additional-options "-fopt-info-optimized-omp" } + { dg-additional-options "-fopt-info-all-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccloops" } */ @@ -23,6 +25,7 @@ extern unsigned int f (unsigned int); void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (unsigned int i = 0; i < N; i++) c[i] = a[f (i)] + b[f (i)]; } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index bdf7b4a0641..b54a71e788a 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -1,8 +1,10 @@ /* Check offloaded function's attributes and classification for OpenACC kernels. */ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + /* { dg-additional-options "-O2" } - { dg-additional-options "-fopt-info-optimized-omp" } + { dg-additional-options "-fopt-info-all-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccloops" } */ @@ -19,6 +21,7 @@ extern unsigned int *__restrict c; void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c index 4dd55eb4680..64ce8943a02 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -55,7 +55,7 @@ main () ; } - { /*TODO Instead of using 'for (int i = 0; [...])', move 'int i' outside, to work around for ICE detailed in 'kernels-decompose-ice-1.c'. */ + { int i; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */ @@ -64,6 +64,20 @@ main () a[i] = 0; } +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ + { + int i; + } + +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + for (int i = 0; i < N; i++) + a[i] = 0; + #pragma acc kernels loop /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c deleted file mode 100644 index e83b451f2b8..00000000000 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c +++ /dev/null @@ -1,114 +0,0 @@ -/* Test OpenACC 'kernels' construct decomposition. */ - -/* { dg-additional-options "-fopt-info-omp-all" } */ - -/* { dg-additional-options "-fchecking --param=openacc-kernels=decompose" } */ -/* { dg-ice "TODO" } - { dg-prune-output "during GIMPLE pass: omplower" } */ - -/* { dg-additional-options "--param=openacc-privatization=noisy" } */ - -/* Reduced from 'kernels-decompose-2.c'. - (Hopefully) similar instances: - - 'kernels-decompose-ice-2.c' - - 'libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c' - - 'libgomp.oacc-c-c++-common/kernels-decompose-1.c' -*/ - -int -main () -{ -#define N 10 - -#pragma acc kernels - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ - /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ - for (int i = 0; i < N; i++) - ; - - return 0; -} - -/* - In 'gimple' we've got: - - main () - { - int D.2087; - - { - int a[10]; - - try - { - #pragma omp target oacc_kernels map(tofrom:a [len: 40]) - { - { - int i; - - i = 0; - goto <D.2085>; - [...] - - ..., which in 'omp_oacc_kernels_decompose' we turn into: - - main () - { - int D.2087; - - { - int a[10]; - - try - { - #pragma omp target oacc_data_kernels map(tofrom:a [len: 40]) - { - try - { - { - int i; - - #pragma omp target oacc_data_kernels map(alloc:i [len: 4]) - { - try - { - { - #pragma omp target oacc_kernels async(-1) map(force_present:i [len: 4]) map(force_present:a [len: 40]) - { - i = 0; - goto <D.2085>; - [...] - - ..., which results in ICE in: - - #1 0x0000000000d2247b in lower_omp_target (gsi_p=gsi_p@entry=0x7fffffffbc90, ctx=ctx@entry=0x2c994c0) at [...]/gcc/omp-low.c:11981 - 11981 gcc_assert (offloaded); - (gdb) list - 11976 talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar))); - 11977 gimplify_assign (x, var, &ilist); - 11978 } - 11979 else if (is_gimple_reg (var)) - 11980 { - 11981 gcc_assert (offloaded); - 11982 tree avar = create_tmp_var (TREE_TYPE (var)); - 11983 mark_addressable (avar); - 11984 enum gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (c); - 11985 if (GOMP_MAP_COPY_TO_P (map_kind) - (gdb) call debug_tree(var) - <var_decl 0x7ffff7feebd0 i - type <integer_type 0x7ffff67be5e8 int sizes-gimplified public SI - size <integer_cst 0x7ffff67a5f18 constant 32> - unit-size <integer_cst 0x7ffff67a5f30 constant 4> - align:32 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff67be5e8 precision:32 min <integer_cst 0x7ffff67a5ed0 -2147483648> max <integer_cst 0x7ffff67a5ee8 2147483647> - pointer_to_this <pointer_type 0x7ffff67c69d8>> - used read SI [...]:15:12 size <integer_cst 0x7ffff67a5f18 32> unit-size <integer_cst 0x7ffff67a5f30 4> - align:32 warn_if_not_align:0 context <function_decl 0x7ffff68eea00 main>> - - Just defusing the 'assert' is not sufficient: - - libgomp: present clause: !acc_is_present (0x7ffe29cba3ec, 4 (0x4)) - - TODO Can't the 'omp_oacc_kernels_decompose' transformation be much simpler, such that we avoid the intermediate 'data' if we've got just one compute construct inside it? - TODO But it's not clear if that'd just resolve one simple instance of the general problem? - -*/ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c deleted file mode 100644 index 16af57d5f87..00000000000 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c +++ /dev/null @@ -1,22 +0,0 @@ -/* Test OpenACC 'kernels' construct decomposition. */ - -/* { dg-additional-options "-fopt-info-omp-all" } */ - -/* { dg-additional-options "-fchecking --param=openacc-kernels=decompose" } */ -/* { dg-ice "TODO" } - { dg-prune-output "during GIMPLE pass: omplower" } */ - -/* { dg-additional-options "--param=openacc-privatization=noisy" } */ - -/* Reduced from 'kernels-decompose-ice-1.c'. */ - -int -main () -{ -#pragma acc kernels - /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .-1 } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ - { - int i; - } -} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c new file mode 100644 index 00000000000..b497af298bc --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c @@ -0,0 +1,19 @@ +/* Reduced from 'libgomp.oacc-c-c++-common/kernels-loop-2.c'. */ + +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +void +foo (void) /* { dg-line l_f_1 } */ +{ +#pragma acc kernels /* { dg-line l_k_1 } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_k_1 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_k_1 } */ + /* { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {TODO location} { xfail *-*-* } l_f_1 } + { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} TODO { xfail *-*-* } .+1 } */ + for (int i;;) + ; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-parloops.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-parloops.f95 new file mode 100644 index 00000000000..b8c2d99a8a4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-parloops.f95 @@ -0,0 +1,43 @@ +! Check offloaded function's attributes and classification for OpenACC +! kernels. + +! { dg-additional-options "--param openacc-kernels=parloops" } + +! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } +! { dg-additional-options "-fdump-tree-ompexp" } +! { dg-additional-options "-fdump-tree-parloops1-all" } +! { dg-additional-options "-fdump-tree-oaccloops" } + +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + +program main + implicit none + integer, parameter :: n = 1024 + integer, dimension (0:n-1) :: a, b, c + integer :: i + + call setup(a, b) + + !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + do i = 0, n - 1 + c(i) = a(i) + b(i) + end do + !$acc end kernels +end program main + +! Check the offloaded function's attributes. +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels, omp target entrypoint\\)\\)" 1 "ompexp" } } + +! Check that exactly one OpenACC kernels construct is analyzed, and that it +! can be parallelized. +! { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops1" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "parloops1" } } +! { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } + +! Check the offloaded function's classification and compute dimensions (will +! always be 1 x 1 x 1 for non-offloading compilation). +! { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized-parloops.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized-parloops.f95 new file mode 100644 index 00000000000..37733273440 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized-parloops.f95 @@ -0,0 +1,47 @@ +! Check offloaded function's attributes and classification for unparallelized +! OpenACC kernels. + +! { dg-additional-options "--param openacc-kernels=parloops" } + +! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } +! { dg-additional-options "-fdump-tree-ompexp" } +! { dg-additional-options "-fdump-tree-parloops1-all" } +! { dg-additional-options "-fdump-tree-oaccloops" } + +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + +program main + implicit none + integer, parameter :: n = 1024 + integer, dimension (0:n-1) :: a, b, c + integer :: i + + ! An "external" mapping of loop iterations/array indices makes the loop + ! unparallelizable. + integer, external :: f + + call setup(a, b) + + !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + do i = 0, n - 1 + c(i) = a(f (i)) + b(f (i)) + end do + !$acc end kernels +end program main + +! Check the offloaded function's attributes. +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels, omp target entrypoint\\)\\)" 1 "ompexp" } } + +! Check that exactly one OpenACC kernels construct is analyzed, and that it +! can't be parallelized. +! { dg-final { scan-tree-dump-times "FAILED:" 1 "parloops1" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "parloops1" } } +! { dg-final { scan-tree-dump-not "SUCCESS: may be parallelized" "parloops1" } } + +! Check the offloaded function's classification and compute dimensions (will +! always be 1 x 1 x 1 for non-offloading compilation). +! { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index 3fb48b321f2..ee8e2899a8a 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -1,8 +1,10 @@ ! Check offloaded function's attributes and classification for unparallelized ! OpenACC kernels. +! { dg-additional-options "--param openacc-kernels=decompose" } + ! { dg-additional-options "-O2" } -! { dg-additional-options "-fopt-info-optimized-omp" } +! { dg-additional-options "-fopt-info-all-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccloops" } @@ -23,6 +25,7 @@ program main call setup(a, b) !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 0, n - 1 c(i) = a(f (i)) + b(f (i)) end do diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index 6c8d298e236..a4bcca05cc1 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -1,8 +1,10 @@ ! Check offloaded function's attributes and classification for OpenACC ! kernels. +! { dg-additional-options "--param openacc-kernels=decompose" } + ! { dg-additional-options "-O2" } -! { dg-additional-options "-fopt-info-optimized-omp" } +! { dg-additional-options "-fopt-info-all-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccloops" } @@ -19,6 +21,7 @@ program main call setup(a, b) !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 0, n - 1 c(i) = a(i) + b(i) end do diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c index a6eb82b8719..3e5b6bab233 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c @@ -1,5 +1,5 @@ /* { dg-additional-options "--param=openacc-kernels=decompose" } */ -/* Hopefully, this is the same issue as '../../../gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c'. +/* ICE similar to PR100280, but not the same. { dg-ice "TODO" } TODO { dg-prune-output "during GIMPLE pass: omplower" } TODO { dg-do link } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c index e4e58158cf7..f7ccecbf4b4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c @@ -3,7 +3,7 @@ /* Based on '../libgomp.oacc-fortran/asyncwait-1.f90'. */ /* { dg-additional-options "--param=openacc-kernels=decompose" } */ -/* TODO To avoid PR100280 ICE { dg-additional-options "--param=openacc-kernels=parloops" } */ +/* { dg-xfail-run-if TODO { openacc_radeon_accel_selected } } */ /* { dg-additional-options "-fopt-info-all-omp" } { dg-additional-options "-foffload=-fopt-info-all-omp" } */ @@ -202,11 +202,12 @@ main (void) #pragma acc data copy (a[0:N]) copy (b[0:N]) { -#pragma acc kernels async /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ - /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } */ +#pragma acc kernels async #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */ for (int i = 0; i < N; ++i) b[i] = a[i]; @@ -229,11 +230,12 @@ main (void) #pragma acc data copy (a[0:N]) copy (b[0:N]) { -#pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ - /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } */ +#pragma acc kernels async (1) #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */ for (int i = 0; i < N; ++i) b[i] = a[i]; @@ -259,24 +261,27 @@ main (void) #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (int i = 0; i < N; ++i) b[i] = (a[i] * a[i] * a[i]) / a[i]; #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (int i = 0; i < N; ++i) c[i] = (a[i] * 4) / a[i]; -#pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ - /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } */ +#pragma acc kernels async (1) #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */ for (int i = 0; i < N; ++i) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; @@ -307,33 +312,37 @@ main (void) #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (int i = 0; i < N; ++i) b[i] = (a[i] * a[i] * a[i]) / a[i]; -#pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ - /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } */ +#pragma acc kernels async (1) #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */ for (int i = 0; i < N; ++i) c[i] = (a[i] * 4) / a[i]; -#pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ - /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } */ +#pragma acc kernels async (1) #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */ for (int i = 0; i < N; ++i) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; -#pragma acc kernels wait (1) async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ - /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } */ +#pragma acc kernels wait (1) async (1) #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */ for (int i = 0; i < N; ++i) e[i] = a[i] + b[i] + c[i] + d[i]; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index b3b4c490f7f..57e75f6d399 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -32,11 +32,7 @@ int main() { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ int c = 234; - /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } - { dg-note {variable 'c\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */ - - /*TODO Hopefully, this is the same issue as '../../../gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c'. */ - (volatile int *) &c; + /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ #pragma acc loop independent gang /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_i$c_loop_i } */ -- 2.34.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] 2022-01-13 9:54 ` OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] Thomas Schwinge @ 2022-03-01 16:46 ` Thomas Schwinge 2022-03-02 13:30 ` Jakub Jelinek ` (3 more replies) 2022-03-04 13:37 ` Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280] Thomas Schwinge 2022-03-04 13:41 ` OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280] Thomas Schwinge 2 siblings, 4 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-01 16:46 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches, fortran; +Cc: asolokha, Julian Brown Hi! Jakub, need your review/approval here, please: On 2022-01-13T10:54:16+0100, I wrote: > On 2019-05-08T14:51:57+0100, Julian Brown <julian@codesourcery.com> wrote: >> - The "addressable" bit is set during the kernels conversion pass for >> variables that have "create" (alloc) clauses created for them in the >> synthesised outer data region (instead of in the front-end, etc., >> where it can't be done accurately). Such variables actually have >> their address taken during transformations made in a later pass >> (omp-low, I think), but there's a phase-ordering problem that means >> the flag should be set earlier. > > The actual issue is a bit different, but yes, there is a problem. > The related ICE has also been reported as <https://gcc.gnu.org/PR100280> > "ICE in lower_omp_target, at omp-low.c:12287". (And I'm confused why we > didn't run into that with the OpenACC 'kernels' decomposition > originally.) I've pushed to master branch > commit 9b32c1669aad5459dd053424f9967011348add83 > "OpenACC 'kernels' decomposition: Mark variables used in synthesized data > clauses as addressable [PR100280]" > ... as otherwise 'gcc/omp-low.c:lower_omp_target' has to create a temporary: > > 13073 else if (is_gimple_reg (var)) > 13074 { > 13075 gcc_assert (offloaded); > 13076 tree avar = create_tmp_var (TREE_TYPE (var)); > 13077 mark_addressable (avar); > > ..., which (a) is only implemented for actualy *offloaded* regions (but not > data regions), and (b) the subsequently synthesized code for writing to and > later reading back from the temporary fundamentally conflicts with OpenACC > 'async' (as used by OpenACC 'kernels' decomposition). That's all not trivial > to make work, so let's just avoid this case. > --- a/gcc/omp-oacc-kernels-decompose.cc > +++ b/gcc/omp-oacc-kernels-decompose.cc > @@ -793,7 +793,8 @@ make_data_region_try_statement (location_t loc, gimple *body) > > /* If INNER_BIND_VARS holds variables, build an OpenACC data region with > location LOC containing BODY and having 'create (var)' clauses for each > - variable. If INNER_CLEANUP is present, add a try-finally statement with > + variable (as a side effect, such variables also get TREE_ADDRESSABLE set). > + If INNER_CLEANUP is present, add a try-finally statement with > this cleanup code in the finally block. Return the new data region, or > the original BODY if no data region was needed. */ > > @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, > inner_data_clauses = new_clause; > > prev_mapped_var = v; > + > + /* See <https://gcc.gnu.org/PR100280>. */ > + TREE_ADDRESSABLE (v) = 1; > } > } So, that's too simple. ;-) ... and gives rise to workaround patches like we have on the og11 development branch: - "Avoid introducing 'create' mapping clauses for loop index variables in kernels regions", - "Run all kernels regions with GOMP_MAP_FORCE_TOFROM mappings synchronously", - "Fix for is_gimple_reg vars to 'data kernels'" We're after gimplification, and must not just set 'TREE_ADDRESSABLE', because that may easily violate GIMPLE invariants, leading to ICEs later. There are a few open PRs, which my following changes are addressing. To make "late" 'TREE_ADDRESSABLE' work, we have a precedent in OpenMP's 'gcc/omp-low.cc:task_shared_vars' handling, as Jakub had pointed to in discussion of <https://gcc.gnu.org/PR102330>. (PR102330 turned out to be unrelated from the "late" 'TREE_ADDRESSABLE' problem here; I have a different patch for it.) I'm thus proposing to generalize 'gcc/omp-low.cc:task_shared_vars' into 'make_addressable_vars', plus new 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' that we then may use instead of the 'TREE_ADDRESSABLE (v) = 1;' quoted above (plus one or two additional ones to be introduced in later patches), and wire that up in 'gcc/omp-low.cc:scan_sharing_clauses', for 'OMP_CLAUSE_MAP': set 'TREE_ADDRESSABLE' and put into 'make_addressable_vars' for later fix-up. (In reply to Jakub Jelinek from comment #9) > Whether you can use the same bitmap or need to add another bitmap next to > task_shared_vars is something hard to guess without diving into it deeply. Per my understanding of the code, the only place where I had doubts is 'gcc/omp-low.cc:finish_taskreg_scan', but I have convinced myself that what this is doing is either a no-op in the 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' case, or in fact necessary as the original 'task_shared_vars' handling has been. Either way: I couldn't come up with a way (test case) that we'd actually run into this case; you'd have to have the relevant OpenMP constructs inside an OpenACC 'kernels' region, which isn't permitted per 'gcc/omp-low.cc:check_omp_nesting_restrictions'. OK to proceed in this way? Grüße Thomas --- gcc/omp-low.cc +++ gcc/omp-low.cc @@ -188,7 +188,7 @@ struct omp_context static splay_tree all_contexts; static int taskreg_nesting_level; static int target_nesting_level; -static bitmap task_shared_vars; +static bitmap make_addressable_vars; static bitmap global_nonaddressable_vars; static vec<omp_context *> taskreg_contexts; static vec<gomp_task *> task_cpyfns; @@ -572,9 +572,9 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx) /* Taking address of OUTER in lower_send_shared_vars might need regimplification of everything that uses the variable. */ - if (!task_shared_vars) - task_shared_vars = BITMAP_ALLOC (NULL); - bitmap_set_bit (task_shared_vars, DECL_UID (outer)); + if (!make_addressable_vars) + make_addressable_vars = BITMAP_ALLOC (NULL); + bitmap_set_bit (make_addressable_vars, DECL_UID (outer)); TREE_ADDRESSABLE (outer) = 1; } return true; @@ -601,13 +601,13 @@ omp_copy_decl_2 (tree var, tree name, tree type, omp_context *ctx) else record_vars (copy); - /* If VAR is listed in task_shared_vars, it means it wasn't - originally addressable and is just because task needs to take - it's address. But we don't need to take address of privatizations + /* If VAR is listed in make_addressable_vars, it wasn't + originally addressable, but was only later made so. + We don't need to take address of privatizations from that var. */ if (TREE_ADDRESSABLE (var) - && ((task_shared_vars - && bitmap_bit_p (task_shared_vars, DECL_UID (var))) + && ((make_addressable_vars + && bitmap_bit_p (make_addressable_vars, DECL_UID (var))) || (global_nonaddressable_vars && bitmap_bit_p (global_nonaddressable_vars, DECL_UID (var))))) TREE_ADDRESSABLE (copy) = 0; @@ -1495,6 +1495,21 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); decl = OMP_CLAUSE_DECL (c); + /* If requested, make 'decl' addressable. */ + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c)) + { + gcc_checking_assert (DECL_P (decl)); + + gcc_checking_assert (!TREE_ADDRESSABLE (decl)); + if (!make_addressable_vars) + make_addressable_vars = BITMAP_ALLOC (NULL); + bitmap_set_bit (make_addressable_vars, DECL_UID (decl)); + TREE_ADDRESSABLE (decl) = 1; + + /* Done. */ + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c) = 0; + } /* Global variables with "omp declare target" attribute don't need to be copied, the receiver side will use them directly. However, global variables with "omp declare target link" @@ -2371,11 +2405,11 @@ finish_taskreg_scan (omp_context *ctx) if (ctx->record_type == NULL_TREE) return; - /* If any task_shared_vars were needed, verify all + /* If any make_addressable_vars were needed, verify all OMP_CLAUSE_SHARED clauses on GIMPLE_OMP_{PARALLEL,TASK,TEAMS} statements if use_pointer_for_field hasn't changed because of that. If it did, update field types now. */ - if (task_shared_vars) + if (make_addressable_vars) { tree c; @@ -2390,7 +2424,7 @@ finish_taskreg_scan (omp_context *ctx) the receiver side will use them directly. */ if (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))) continue; - if (!bitmap_bit_p (task_shared_vars, DECL_UID (decl)) + if (!bitmap_bit_p (make_addressable_vars, DECL_UID (decl)) || !use_pointer_for_field (decl, ctx)) continue; tree field = lookup_field (decl, ctx); @@ -14040,7 +14074,7 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx) /* Callback for lower_omp_1. Return non-NULL if *tp needs to be regimplified. If DATA is non-NULL, lower_omp_1 is outside - of OMP context, but with task_shared_vars set. */ + of OMP context, but with make_addressable_vars set. */ static tree lower_omp_regimplify_p (tree *tp, int *walk_subtrees, @@ -14054,9 +14088,9 @@ lower_omp_regimplify_p (tree *tp, int *walk_subtrees, && DECL_HAS_VALUE_EXPR_P (t)) return t; - if (task_shared_vars + if (make_addressable_vars && DECL_P (t) - && bitmap_bit_p (task_shared_vars, DECL_UID (t))) + && bitmap_bit_p (make_addressable_vars, DECL_UID (t))) return t; /* If a global variable has been privatized, TREE_CONSTANT on @@ -14141,7 +14175,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (gimple_has_location (stmt)) input_location = gimple_location (stmt); - if (task_shared_vars) + if (make_addressable_vars) memset (&wi, '\0', sizeof (wi)); /* If we have issued syntax errors, avoid doing any heavy lifting. @@ -14158,7 +14192,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GIMPLE_COND: { gcond *cond_stmt = as_a <gcond *> (stmt); - if ((ctx || task_shared_vars) + if ((ctx || make_addressable_vars) && (walk_tree (gimple_cond_lhs_ptr (cond_stmt), lower_omp_regimplify_p, ctx ? NULL : &wi, NULL) @@ -14250,7 +14284,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) lower_omp_critical (gsi_p, ctx); break; case GIMPLE_OMP_ATOMIC_LOAD: - if ((ctx || task_shared_vars) + if ((ctx || make_addressable_vars) && walk_tree (gimple_omp_atomic_load_rhs_ptr ( as_a <gomp_atomic_load *> (stmt)), lower_omp_regimplify_p, ctx ? NULL : &wi, NULL)) @@ -14371,7 +14405,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: regimplify: - if ((ctx || task_shared_vars) + if ((ctx || make_addressable_vars) && walk_gimple_op (stmt, lower_omp_regimplify_p, ctx ? NULL : &wi)) { @@ -14435,10 +14469,10 @@ execute_lower_omp (void) if (all_contexts->root) { - if (task_shared_vars) + if (make_addressable_vars) push_gimplify_context (); lower_omp (&body, NULL); - if (task_shared_vars) + if (make_addressable_vars) pop_gimplify_context (NULL); } @@ -14447,7 +14481,7 @@ execute_lower_omp (void) splay_tree_delete (all_contexts); all_contexts = NULL; } - BITMAP_FREE (task_shared_vars); + BITMAP_FREE (make_addressable_vars); BITMAP_FREE (global_nonaddressable_vars); /* If current function is a method, remove artificial dummy VAR_DECL created --- gcc/omp-oacc-kernels-decompose.cc +++ gcc/omp-oacc-kernels-decompose.cc @@ -845,7 +845,11 @@ maybe_build_inner_data_region (location_t loc, gimple *body, prev_mapped_var = v; /* See <https://gcc.gnu.org/PR100280>. */ - TREE_ADDRESSABLE (v) = 1; + if (!TREE_ADDRESSABLE (v)) + { + /* Request that OMP lowering make 'v' addressable. */ + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1; + } } } --- gcc/tree-core.h +++ gcc/tree-core.h @@ -1145,6 +1145,9 @@ struct GTY(()) tree_base { PREDICT_EXPR_OUTCOME in PREDICT_EXPR + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE in + OMP_CLAUSE + static_flag: TREE_STATIC in --- gcc/tree.h +++ gcc/tree.h @@ -1695,6 +1695,11 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag) +/* Flag that 'OMP_CLAUSE_DECL (NODE)' is to be made addressable during OMP + lowering. */ +#define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag) + /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' clause. */ #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \ ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] 2022-03-01 16:46 ` Thomas Schwinge @ 2022-03-02 13:30 ` Jakub Jelinek 2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104132-1.c' [PR104132] Thomas Schwinge ` (2 subsequent siblings) 3 siblings, 0 replies; 10+ messages in thread From: Jakub Jelinek @ 2022-03-02 13:30 UTC (permalink / raw) To: Thomas Schwinge; +Cc: gcc-patches, fortran, asolokha, Julian Brown On Tue, Mar 01, 2022 at 05:46:20PM +0100, Thomas Schwinge wrote: > OK to proceed in this way? With a suitable ChangeLog entry and one nit fixed yes. > --- gcc/omp-low.cc > +++ gcc/omp-low.cc > @@ -188,7 +188,7 @@ struct omp_context > static splay_tree all_contexts; > static int taskreg_nesting_level; > static int target_nesting_level; > -static bitmap task_shared_vars; > +static bitmap make_addressable_vars; > static bitmap global_nonaddressable_vars; > static vec<omp_context *> taskreg_contexts; > static vec<gomp_task *> task_cpyfns; > @@ -572,9 +572,9 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx) > /* Taking address of OUTER in lower_send_shared_vars > might need regimplification of everything that uses the > variable. */ > - if (!task_shared_vars) > - task_shared_vars = BITMAP_ALLOC (NULL); > - bitmap_set_bit (task_shared_vars, DECL_UID (outer)); > + if (!make_addressable_vars) > + make_addressable_vars = BITMAP_ALLOC (NULL); > + bitmap_set_bit (make_addressable_vars, DECL_UID (outer)); Has the MUA replaced tabs with spaces? > --- gcc/omp-oacc-kernels-decompose.cc > +++ gcc/omp-oacc-kernels-decompose.cc > @@ -845,7 +845,11 @@ maybe_build_inner_data_region (location_t loc, gimple *body, > prev_mapped_var = v; > > /* See <https://gcc.gnu.org/PR100280>. */ > - TREE_ADDRESSABLE (v) = 1; > + if (!TREE_ADDRESSABLE (v)) > + { > + /* Request that OMP lowering make 'v' addressable. */ > + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1; > + } That is a single statement body, so shouldn't have {}s around it. Jakub ^ permalink raw reply [flat|nested] 10+ messages in thread
* Add 'c-c++-common/goacc/kernels-decompose-pr104132-1.c' [PR104132] 2022-03-01 16:46 ` Thomas Schwinge 2022-03-02 13:30 ` Jakub Jelinek @ 2022-03-04 13:34 ` Thomas Schwinge 2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104133-1.c' [PR104133] Thomas Schwinge 2022-03-04 13:46 ` OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] Thomas Schwinge 3 siblings, 0 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-04 13:34 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek, Julian Brown, asolokha, Andrew Pinski [-- Attachment #1: Type: text/plain, Size: 1143 bytes --] Hi! On 2022-03-01T17:46:20+0100, I wrote: > On 2022-01-13T10:54:16+0100, I wrote: >> --- a/gcc/omp-oacc-kernels-decompose.cc >> +++ b/gcc/omp-oacc-kernels-decompose.cc >> @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, >> inner_data_clauses = new_clause; >> >> prev_mapped_var = v; >> + >> + /* See <https://gcc.gnu.org/PR100280>. */ >> + TREE_ADDRESSABLE (v) = 1; >> } >> } > > So, that's too simple. ;-) [...] > We're after gimplification, and must not just set 'TREE_ADDRESSABLE', > because that may easily violate GIMPLE invariants, leading to ICEs later. > There are a few open PRs Pushed to master branch commit 741859b390c042755e9379f8061a157e5af378b6 "Add 'c-c++-common/goacc/kernels-decompose-pr104132-1.c' [PR104132]", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Add-c-c-common-goacc-kernels-decompose-pr104132-1.c-.patch --] [-- Type: text/x-diff, Size: 2619 bytes --] From 741859b390c042755e9379f8061a157e5af378b6 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Wed, 19 Jan 2022 22:28:55 +0100 Subject: [PATCH] Add 'c-c++-common/goacc/kernels-decompose-pr104132-1.c' [PR104132] ..., currently XFAILed with 'dg-ice'. PR middle-end/104132 gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-pr104132-1.c: New file. --- .../goacc/kernels-decompose-pr104132-1.c | 38 +++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c new file mode 100644 index 00000000000..4fbfdd81e15 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c @@ -0,0 +1,38 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO } + { dg-prune-output {k = 0 \+ \.offset\.[0-9]+;} } + { dg-prune-output {k = 0 \+ 2;} } + { dg-prune-output {during IPA pass: \*free_lang_data} } */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } + Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ + +int arr_0; + +void +foo (void) +{ +#pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ + { + int k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_k1 } */ + /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k1 } */ + for (k = 0; k < 2; k++) + arr_0 = k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_k2 } */ + /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k2 } */ + for (k = 0; k < 2; k++) + arr_0 = k; + } +} +/* { dg-bogus {error: non-register as LHS of binary operation} {} { xfail *-*-* } .-1 } */ -- 2.34.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Add 'c-c++-common/goacc/kernels-decompose-pr104133-1.c' [PR104133] 2022-03-01 16:46 ` Thomas Schwinge 2022-03-02 13:30 ` Jakub Jelinek 2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104132-1.c' [PR104132] Thomas Schwinge @ 2022-03-04 13:34 ` Thomas Schwinge 2022-03-04 13:46 ` OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] Thomas Schwinge 3 siblings, 0 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-04 13:34 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek, Julian Brown, asolokha, Andrew Pinski [-- Attachment #1: Type: text/plain, Size: 1143 bytes --] Hi! On 2022-03-01T17:46:20+0100, I wrote: > On 2022-01-13T10:54:16+0100, I wrote: >> --- a/gcc/omp-oacc-kernels-decompose.cc >> +++ b/gcc/omp-oacc-kernels-decompose.cc >> @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, >> inner_data_clauses = new_clause; >> >> prev_mapped_var = v; >> + >> + /* See <https://gcc.gnu.org/PR100280>. */ >> + TREE_ADDRESSABLE (v) = 1; >> } >> } > > So, that's too simple. ;-) [...] > We're after gimplification, and must not just set 'TREE_ADDRESSABLE', > because that may easily violate GIMPLE invariants, leading to ICEs later. > There are a few open PRs Pushed to master branch commit e085900fa10e28b684d656b66557d181247a1a48 "Add 'c-c++-common/goacc/kernels-decompose-pr104133-1.c' [PR104133]", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Add-c-c-common-goacc-kernels-decompose-pr104133-1.c-.patch --] [-- Type: text/x-diff, Size: 2965 bytes --] From e085900fa10e28b684d656b66557d181247a1a48 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Wed, 19 Jan 2022 22:28:55 +0100 Subject: [PATCH] Add 'c-c++-common/goacc/kernels-decompose-pr104133-1.c' [PR104133] ..., currently XFAILed with 'dg-ice'. PR middle-end/104133 gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-pr104133-1.c: New file. --- .../goacc/kernels-decompose-pr104133-1.c | 40 +++++++++++++++++++ 1 file changed, 40 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c new file mode 100644 index 00000000000..72dde346dbf --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c @@ -0,0 +1,40 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO } + { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} } + { dg-prune-output {D\.[0-9]+ = arr_0\.1 \+ k;} } + { dg-prune-output {during GIMPLE pass: lower} } */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } + Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ + +int arr_0; + +void +foo (void) +{ +#pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'arr_0\.1' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ + { + int k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_k1 } */ + /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k1 } */ + for (k = 0; k < 2; k++) + arr_0 += k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_k2 } */ + /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k2 } */ + for (k = 0; k < 2; k++) + arr_0 += k; + /* { dg-bogus {error: invalid operands in binary operation} {} { xfail *-*-* } .-1 } */ + } +} -- 2.34.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
* OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] 2022-03-01 16:46 ` Thomas Schwinge ` (2 preceding siblings ...) 2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104133-1.c' [PR104133] Thomas Schwinge @ 2022-03-04 13:46 ` Thomas Schwinge 2022-03-04 14:57 ` Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784] Thomas Schwinge 2022-03-04 19:52 ` Thomas Schwinge 3 siblings, 2 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-04 13:46 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek, Julian Brown, asolokha, Andrew Pinski [-- Attachment #1: Type: text/plain, Size: 3596 bytes --] Hi! On 2022-03-01T17:46:20+0100, I wrote: > On 2022-01-13T10:54:16+0100, I wrote: >> On 2019-05-08T14:51:57+0100, Julian Brown <julian@codesourcery.com> wrote: >>> - The "addressable" bit is set during the kernels conversion pass for >>> variables that have "create" (alloc) clauses created for them in the >>> synthesised outer data region (instead of in the front-end, etc., >>> where it can't be done accurately). Such variables actually have >>> their address taken during transformations made in a later pass >>> (omp-low, I think), but there's a phase-ordering problem that means >>> the flag should be set earlier. >> >> The actual issue is a bit different, but yes, there is a problem. >> The related ICE has also been reported as <https://gcc.gnu.org/PR100280> >> "ICE in lower_omp_target, at omp-low.c:12287". (And I'm confused why we >> didn't run into that with the OpenACC 'kernels' decomposition >> originally.) I've pushed to master branch >> commit 9b32c1669aad5459dd053424f9967011348add83 >> "OpenACC 'kernels' decomposition: Mark variables used in synthesized data >> clauses as addressable [PR100280]" >> --- a/gcc/omp-oacc-kernels-decompose.cc >> +++ b/gcc/omp-oacc-kernels-decompose.cc >> @@ -793,7 +793,8 @@ make_data_region_try_statement (location_t loc, gimple *body) >> >> /* If INNER_BIND_VARS holds variables, build an OpenACC data region with >> location LOC containing BODY and having 'create (var)' clauses for each >> - variable. If INNER_CLEANUP is present, add a try-finally statement with >> + variable (as a side effect, such variables also get TREE_ADDRESSABLE set). >> + If INNER_CLEANUP is present, add a try-finally statement with >> this cleanup code in the finally block. Return the new data region, or >> the original BODY if no data region was needed. */ >> >> @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, >> inner_data_clauses = new_clause; >> >> prev_mapped_var = v; >> + >> + /* See <https://gcc.gnu.org/PR100280>. */ >> + TREE_ADDRESSABLE (v) = 1; >> } >> } > > So, that's too simple. ;-) [...] > We're after gimplification, and must not just set 'TREE_ADDRESSABLE', > because that may easily violate GIMPLE invariants, leading to ICEs later. > There are a few open PRs, which my following changes are addressing. To > make "late" 'TREE_ADDRESSABLE' work, we have a precedent in OpenMP's > 'gcc/omp-low.cc:task_shared_vars' handling, as Jakub had pointed to in > discussion of <https://gcc.gnu.org/PR102330>. > I'm thus proposing to generalize 'gcc/omp-low.cc:task_shared_vars' into > 'make_addressable_vars', plus new 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' > that we then may use instead of the 'TREE_ADDRESSABLE (v) = 1;' quoted > above (plus one or two additional ones to be introduced in later > patches), and wire that up in 'gcc/omp-low.cc:scan_sharing_clauses', for > 'OMP_CLAUSE_MAP': set 'TREE_ADDRESSABLE' and put into > 'make_addressable_vars' for later fix-up. Pushed to master branch commit 8935589b496f755e08cadf26d8ceddf0dd6e0968 "OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133]", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-OMP-lowering-Regimplify-OMP_CLAUSE_MAP_DECL_MAKE_ADD.patch --] [-- Type: text/x-diff, Size: 20858 bytes --] From 8935589b496f755e08cadf26d8ceddf0dd6e0968 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Tue, 15 Feb 2022 23:31:34 +0100 Subject: [PATCH] OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] ... by generalizing the existing 'gcc/omp-low.cc:task_shared_vars'. Fix-up for commit 9b32c1669aad5459dd053424f9967011348add83 "OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280]". PR middle-end/100280 PR middle-end/104132 PR middle-end/104133 gcc/ * omp-low.cc (task_shared_vars): Rename to 'make_addressable_vars'. Adjust all users. (scan_sharing_clauses) <OMP_CLAUSE_MAP> Use it for 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs, too. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Adjust. * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Extend. --- gcc/omp-low.cc | 47 +++++++------- .../goacc/kernels-decompose-pr104061-1-3.c | 11 +--- .../goacc/kernels-decompose-pr104061-1-4.c | 17 ++--- .../goacc/kernels-decompose-pr104132-1.c | 9 +-- .../goacc/kernels-decompose-pr104133-1.c | 9 +-- .../kernels-decompose-1.c | 62 +++++++++++++++++-- 6 files changed, 95 insertions(+), 60 deletions(-) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 6654bfd426e..5ce3a50709a 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -188,7 +188,7 @@ struct omp_context static splay_tree all_contexts; static int taskreg_nesting_level; static int target_nesting_level; -static bitmap task_shared_vars; +static bitmap make_addressable_vars; static bitmap global_nonaddressable_vars; static vec<omp_context *> taskreg_contexts; static vec<gomp_task *> task_cpyfns; @@ -572,9 +572,9 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx) /* Taking address of OUTER in lower_send_shared_vars might need regimplification of everything that uses the variable. */ - if (!task_shared_vars) - task_shared_vars = BITMAP_ALLOC (NULL); - bitmap_set_bit (task_shared_vars, DECL_UID (outer)); + if (!make_addressable_vars) + make_addressable_vars = BITMAP_ALLOC (NULL); + bitmap_set_bit (make_addressable_vars, DECL_UID (outer)); TREE_ADDRESSABLE (outer) = 1; } return true; @@ -601,13 +601,13 @@ omp_copy_decl_2 (tree var, tree name, tree type, omp_context *ctx) else record_vars (copy); - /* If VAR is listed in task_shared_vars, it means it wasn't - originally addressable and is just because task needs to take - it's address. But we don't need to take address of privatizations + /* If VAR is listed in make_addressable_vars, it wasn't + originally addressable, but was only later made so. + We don't need to take address of privatizations from that var. */ if (TREE_ADDRESSABLE (var) - && ((task_shared_vars - && bitmap_bit_p (task_shared_vars, DECL_UID (var))) + && ((make_addressable_vars + && bitmap_bit_p (make_addressable_vars, DECL_UID (var))) || (global_nonaddressable_vars && bitmap_bit_p (global_nonaddressable_vars, DECL_UID (var))))) TREE_ADDRESSABLE (copy) = 0; @@ -1502,6 +1502,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) gcc_checking_assert (DECL_P (decl)); gcc_checking_assert (!TREE_ADDRESSABLE (decl)); + if (!make_addressable_vars) + make_addressable_vars = BITMAP_ALLOC (NULL); + bitmap_set_bit (make_addressable_vars, DECL_UID (decl)); TREE_ADDRESSABLE (decl) = 1; if (dump_enabled_p ()) @@ -2402,11 +2405,11 @@ finish_taskreg_scan (omp_context *ctx) if (ctx->record_type == NULL_TREE) return; - /* If any task_shared_vars were needed, verify all + /* If any make_addressable_vars were needed, verify all OMP_CLAUSE_SHARED clauses on GIMPLE_OMP_{PARALLEL,TASK,TEAMS} statements if use_pointer_for_field hasn't changed because of that. If it did, update field types now. */ - if (task_shared_vars) + if (make_addressable_vars) { tree c; @@ -2421,7 +2424,7 @@ finish_taskreg_scan (omp_context *ctx) the receiver side will use them directly. */ if (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))) continue; - if (!bitmap_bit_p (task_shared_vars, DECL_UID (decl)) + if (!bitmap_bit_p (make_addressable_vars, DECL_UID (decl)) || !use_pointer_for_field (decl, ctx)) continue; tree field = lookup_field (decl, ctx); @@ -14071,7 +14074,7 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx) /* Callback for lower_omp_1. Return non-NULL if *tp needs to be regimplified. If DATA is non-NULL, lower_omp_1 is outside - of OMP context, but with task_shared_vars set. */ + of OMP context, but with make_addressable_vars set. */ static tree lower_omp_regimplify_p (tree *tp, int *walk_subtrees, @@ -14085,9 +14088,9 @@ lower_omp_regimplify_p (tree *tp, int *walk_subtrees, && DECL_HAS_VALUE_EXPR_P (t)) return t; - if (task_shared_vars + if (make_addressable_vars && DECL_P (t) - && bitmap_bit_p (task_shared_vars, DECL_UID (t))) + && bitmap_bit_p (make_addressable_vars, DECL_UID (t))) return t; /* If a global variable has been privatized, TREE_CONSTANT on @@ -14172,7 +14175,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (gimple_has_location (stmt)) input_location = gimple_location (stmt); - if (task_shared_vars) + if (make_addressable_vars) memset (&wi, '\0', sizeof (wi)); /* If we have issued syntax errors, avoid doing any heavy lifting. @@ -14189,7 +14192,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GIMPLE_COND: { gcond *cond_stmt = as_a <gcond *> (stmt); - if ((ctx || task_shared_vars) + if ((ctx || make_addressable_vars) && (walk_tree (gimple_cond_lhs_ptr (cond_stmt), lower_omp_regimplify_p, ctx ? NULL : &wi, NULL) @@ -14281,7 +14284,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) lower_omp_critical (gsi_p, ctx); break; case GIMPLE_OMP_ATOMIC_LOAD: - if ((ctx || task_shared_vars) + if ((ctx || make_addressable_vars) && walk_tree (gimple_omp_atomic_load_rhs_ptr ( as_a <gomp_atomic_load *> (stmt)), lower_omp_regimplify_p, ctx ? NULL : &wi, NULL)) @@ -14402,7 +14405,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: regimplify: - if ((ctx || task_shared_vars) + if ((ctx || make_addressable_vars) && walk_gimple_op (stmt, lower_omp_regimplify_p, ctx ? NULL : &wi)) { @@ -14466,10 +14469,10 @@ execute_lower_omp (void) if (all_contexts->root) { - if (task_shared_vars) + if (make_addressable_vars) push_gimplify_context (); lower_omp (&body, NULL); - if (task_shared_vars) + if (make_addressable_vars) pop_gimplify_context (NULL); } @@ -14478,7 +14481,7 @@ execute_lower_omp (void) splay_tree_delete (all_contexts); all_contexts = NULL; } - BITMAP_FREE (task_shared_vars); + BITMAP_FREE (make_addressable_vars); BITMAP_FREE (global_nonaddressable_vars); /* If current function is a method, remove artificial dummy VAR_DECL created diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c index f41dda86122..e106fc32c4f 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c @@ -1,12 +1,6 @@ /* { dg-additional-options "--param openacc-kernels=decompose" } */ -/* { dg-additional-options "-fchecking" } - { dg-ice TODO } - { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} } - { dg-prune-output {during GIMPLE pass: lower} } */ - /* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first. - { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail *-*-* } 0 } { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail *-*-* } 0 } */ /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ @@ -35,11 +29,10 @@ foo (void) /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_k1 } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */ - /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */ - /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */ + /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {w/o debug} { target *-*-* } l_loop_k1 } + { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */ for (k = 0; k < 2; k++) arr_0 += k; - /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */ } } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c index cde95a7b7ac..bedbb0a30eb 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c @@ -1,11 +1,7 @@ /* { dg-additional-options "--param openacc-kernels=decompose" } */ -/* { dg-additional-options "-fchecking" } - { dg-ice TODO } - { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} } - { dg-prune-output {during GIMPLE pass: lower} } */ - -/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */ +/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. + { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail *-*-* } 0 } */ /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ /* { dg-additional-options "-fopt-info-all-omp" } */ @@ -32,12 +28,11 @@ foo (void) /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_k1 } */ - /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */ - /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */ - /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } l_loop_k1 } */ + /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */ + /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {w/o debug} { target *-*-* } l_loop_k1 } + { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */ for (k = 0; k < 2; k++) arr_0 += k; - /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */ } } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c index 4f38a83bb19..42ec4418e40 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c @@ -1,11 +1,5 @@ /* { dg-additional-options "--param openacc-kernels=decompose" } */ -/* { dg-additional-options "-fchecking" } - { dg-ice TODO } - { dg-prune-output {k = 0 \+ \.offset\.[0-9]+;} } - { dg-prune-output {k = 0 \+ 2;} } - { dg-prune-output {during IPA pass: \*free_lang_data} } */ - /* { dg-additional-options "-fopt-info-all-omp" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } @@ -27,14 +21,15 @@ foo (void) /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_k1 } */ /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k1 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */ for (k = 0; k < 2; k++) arr_0 = k; /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_k2 } */ /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k2 } */ for (k = 0; k < 2; k++) arr_0 = k; } } -/* { dg-bogus {error: non-register as LHS of binary operation} {} { xfail *-*-* } .-1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c index 0499665777d..47ea2b92959 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c @@ -1,11 +1,5 @@ /* { dg-additional-options "--param openacc-kernels=decompose" } */ -/* { dg-additional-options "-fchecking" } - { dg-ice TODO } - { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} } - { dg-prune-output {D\.[0-9]+ = arr_0\.1 \+ k;} } - { dg-prune-output {during GIMPLE pass: lower} } */ - /* { dg-additional-options "-fopt-info-all-omp" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } @@ -29,14 +23,15 @@ foo (void) /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_k1 } */ /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k1 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */ for (k = 0; k < 2; k++) arr_0 += k; /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_k2 } */ /* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k2 } */ for (k = 0; k < 2; k++) arr_0 += k; - /* { dg-bogus {error: invalid operands in binary operation} {} { xfail *-*-* } .-1 } */ } } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index 85c39871f94..049b3a44b03 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -7,19 +7,23 @@ /* { dg-additional-options "--param=openacc-privatization=noisy" } { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } - for testing/documenting aspects of that functionality. */ + Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", so to maintain compatibility with earlier Tcl releases, we manually initialize counter variables: - { dg-line l_dummy[variable c_compute 0 c_loop_i 0] } + { dg-line l_dummy[variable c_compute 0 c_loop_c 0 c_loop_i 0] } { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid "WARNING: dg-line var l_dummy defined, but not used". */ #undef NDEBUG #include <assert.h> +static int g1; +static int g2; + int main() { int a = 0; @@ -27,8 +31,12 @@ int main() (volatile int *) &a; #define N 123 int b[N] = { 0 }; + unsigned long long f1; + /*TODO See above. */ + (volatile void *) &f1; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {variable 'g2\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ int c = 234; @@ -46,11 +54,57 @@ int main() /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ a = c; + + /* PR104132, PR104133 */ + { + /* Use the 'kernels'-top-level 'int c' as loop variable. */ + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */ + for (c = 0; c < N / 2; c++) + b[c] -= 10; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */ + for (c = 0; c < N / 2; c++) + g1 = c; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */ + for (c = 0; c <= N; c++) + g2 += c; + + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + f1 = 1; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */ + for (c = 20; c > 0; --c) + f1 *= c; + + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + if (c != 234) + __builtin_abort (); + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + } } - for (int i = 0; i < N; ++i) - assert (b[i] == 234); assert (a == 234); + for (int i = 0; i < N; ++i) + if (i < N / 2) + assert (b[i] == 234 - 10); + else + assert (b[i] == 234); + assert (g1 == N / 2 - 1); + assert (g2 == N * (N + 1) / 2); + assert (f1 == 2432902008176640000ULL); return 0; } -- 2.34.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784] 2022-03-04 13:46 ` OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] Thomas Schwinge @ 2022-03-04 14:57 ` Thomas Schwinge 2022-03-04 19:52 ` Thomas Schwinge 1 sibling, 0 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-04 14:57 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek, Julian Brown, asolokha, Andrew Pinski [-- Attachment #1: Type: text/plain, Size: 732 bytes --] Hi! On 2022-03-04T14:46:25+0100, I wrote: > Pushed to master branch commit 8935589b496f755e08cadf26d8ceddf0dd6e0968 > "OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs > [PR100280, PR104132, PR104133]", see attached. Pushed to master branch commit e28eb86c18ed765dceb3c56471a848e9f0e120ff "Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784]", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Test-libgomp.oacc-kernels-private-vars-with-param-op.patch --] [-- Type: text/x-diff, Size: 104687 bytes --] From e28eb86c18ed765dceb3c56471a848e9f0e120ff Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Wed, 16 Feb 2022 22:24:03 +0100 Subject: [PATCH] Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784] Before recent commit 8935589b496f755e08cadf26d8ceddf0dd6e0968 "OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133]", 'libgomp.oacc-c' testing already worked fine, but 'libgomp.oacc-c++' testing ICEed. Via the commit mentioned, the C++ testing ICEs are now resolved, but the underlying issue remains to be looked into: PR104784 "OpenACC 'kernels' decomposition: C vs. C++ differences". PR middle-end/104784 libgomp/ * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c: Test with '--param=openacc-kernels=decompose'. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90: Likewise. --- .../kernels-private-vars-local-worker-1.c | 23 +++++++++++++++---- .../kernels-private-vars-local-worker-2.c | 20 ++++++++++++---- .../kernels-private-vars-local-worker-3.c | 20 ++++++++++++---- .../kernels-private-vars-local-worker-4.c | 20 ++++++++++++---- .../kernels-private-vars-local-worker-5.c | 20 ++++++++++++---- .../kernels-private-vars-loop-gang-1.c | 11 ++++++--- .../kernels-private-vars-loop-gang-2.c | 11 ++++++--- .../kernels-private-vars-loop-gang-3.c | 11 ++++++--- .../kernels-private-vars-loop-gang-4.c | 10 ++++++-- .../kernels-private-vars-loop-gang-5.c | 11 ++++++--- .../kernels-private-vars-loop-gang-6.c | 11 ++++++--- .../kernels-private-vars-loop-vector-1.c | 20 ++++++++++++---- .../kernels-private-vars-loop-vector-2.c | 20 ++++++++++++---- .../kernels-private-vars-loop-worker-1.c | 13 +++++++---- .../kernels-private-vars-loop-worker-2.c | 20 ++++++++++++---- .../kernels-private-vars-loop-worker-3.c | 23 +++++++++++++++---- .../kernels-private-vars-loop-worker-4.c | 20 ++++++++++++---- .../kernels-private-vars-loop-worker-5.c | 21 +++++++++++++---- .../kernels-private-vars-loop-worker-6.c | 20 ++++++++++++---- .../kernels-private-vars-loop-worker-7.c | 20 ++++++++++++---- .../kernels-private-vars-loop-gang-1.f90 | 9 ++++++-- .../kernels-private-vars-loop-gang-2.f90 | 9 ++++++-- .../kernels-private-vars-loop-gang-3.f90 | 9 ++++++-- .../kernels-private-vars-loop-gang-6.f90 | 9 ++++++-- .../kernels-private-vars-loop-vector-1.f90 | 9 ++++++-- .../kernels-private-vars-loop-vector-2.f90 | 9 ++++++-- .../kernels-private-vars-loop-worker-1.f90 | 8 ++++++- .../kernels-private-vars-loop-worker-2.f90 | 9 ++++++-- .../kernels-private-vars-loop-worker-3.f90 | 9 ++++++-- .../kernels-private-vars-loop-worker-4.f90 | 9 ++++++-- .../kernels-private-vars-loop-worker-5.f90 | 8 ++++++- .../kernels-private-vars-loop-worker-6.f90 | 9 ++++++-- .../kernels-private-vars-loop-worker-7.f90 | 9 ++++++-- 33 files changed, 361 insertions(+), 99 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c index acbeb65273f..43bfaf331ab 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared in a local scope, broadcasting to vector-partitioned mode. Back-to-back worker loops. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,16 +30,24 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) @@ -50,7 +62,8 @@ main (int argc, char* argv[]) } #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) @@ -64,8 +77,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += x * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c index 2558a68eb94..c40c2ab33c5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared in a local scope, broadcasting to vector-partitioned mode. Successive vector loops. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,16 +30,24 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) @@ -56,8 +68,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += x * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c index b2a208c163f..bd04dcc7b02 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared in a local scope, broadcasting to vector-partitioned mode. Aggregate worker variable. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -31,16 +35,24 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) @@ -62,8 +74,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += pt.y * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c index 46c395620fe..4303ab848ea 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared in a local scope, broadcasting to vector-partitioned mode. Addressable worker variable. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -31,16 +35,24 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'pt' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'ptp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ @@ -66,8 +78,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += pt.y * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c index 4b5a15e6ad4..8d0e846ce23 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared in a local scope, broadcasting to vector-partitioned mode. Array worker variable. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,16 +30,24 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) @@ -58,8 +70,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += pt[1] * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c index 4a824941427..470d5e65da4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c @@ -1,10 +1,14 @@ /* Test of gang-private variables declared on loop directive. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,8 +28,9 @@ main (int argc, char* argv[]) for (i = 0; i < 32; i++) arr[i] = i; - #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ + #pragma acc kernels copy(arr) { + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ @@ -34,8 +39,8 @@ main (int argc, char* argv[]) x = i * 2; arr[i] += x; } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) assert (arr[i] == i * 3); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c index 039053f3c86..34513e35f22 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c @@ -1,11 +1,15 @@ /* Test of gang-private variables declared on loop directive, with broadcasting to partitioned workers. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -25,8 +29,9 @@ main (int argc, char* argv[]) for (i = 0; i < 32 * 32; i++) arr[i] = i; - #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ + #pragma acc kernels copy(arr) { + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ @@ -40,8 +45,8 @@ main (int argc, char* argv[]) for (int j = 0; j < 32; j++) arr[i * 32 + j] += x; } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32 * 32; i++) assert (arr[i] == i + (i / 32) * 2); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c index 2b89659a007..ef445c82a93 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c @@ -1,11 +1,15 @@ /* Test of gang-private variables declared on loop directive, with broadcasting to partitioned vectors. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -25,8 +29,9 @@ main (int argc, char* argv[]) for (i = 0; i < 32 * 32; i++) arr[i] = i; - #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ + #pragma acc kernels copy(arr) { + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ @@ -40,8 +45,8 @@ main (int argc, char* argv[]) for (int j = 0; j < 32; j++) arr[i * 32 + j] += x; } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32 * 32; i++) assert (arr[i] == i + (i / 32) * 2); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c index 70760705925..0658251ea2a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c @@ -1,11 +1,15 @@ /* Test of gang-private addressable variable declared on loop directive, with broadcasting to partitioned workers. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,7 +30,9 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ @@ -45,8 +51,8 @@ main (int argc, char* argv[]) (*p)--; } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32 * 32; i++) assert (arr[i] == i + (i / 32) * 2); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c index edf0e24af8b..a5118666c63 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c @@ -1,11 +1,15 @@ /* Test of gang-private array variable declared on loop directive, with broadcasting to partitioned workers. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -25,8 +29,9 @@ main (int argc, char* argv[]) for (i = 0; i < 32 * 32; i++) arr[i] = i; - #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ + #pragma acc kernels copy(arr) { + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ @@ -41,8 +46,8 @@ main (int argc, char* argv[]) for (int j = 0; j < 32; j++) arr[i * 32 + j] += x[j % 8]; } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32 * 32; i++) assert (arr[i] == i + (i % 8) * 2); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c index a2df33b767d..119bda3b2dc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c @@ -1,11 +1,15 @@ /* Test of gang-private aggregate variable declared on loop directive, with broadcasting to partitioned workers. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -31,8 +35,9 @@ main (int argc, char* argv[]) for (i = 0; i < 32 * 32; i++) arr[i] = i; - #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ + #pragma acc kernels copy(arr) { + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang private(pt) /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ @@ -49,8 +54,8 @@ main (int argc, char* argv[]) for (int j = 0; j < 32; j++) arr[i * 32 + j] += pt.x + pt.y + pt.z + pt.attr[5]; } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32 * 32; i++) assert (arr[i] == i + (i / 32) * 13); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c index 51c1de53414..5a70bb880e5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c @@ -1,10 +1,14 @@ /* Test of vector-private variables declared on loop directive. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -25,16 +29,24 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -59,8 +71,8 @@ main (int argc, char* argv[]) } } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c index cb90eaab99d..f5bccabbe6f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c @@ -1,10 +1,14 @@ /* Test of vector-private variables declared on loop directive. Array type. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -25,16 +29,24 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c} l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -52,8 +64,8 @@ main (int argc, char* argv[]) } } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c index 54e1c93714b..9b2c49ac301 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c @@ -1,10 +1,14 @@ /* Test of worker-private variables declared on a loop directive. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,12 +28,13 @@ main (int argc, char* argv[]) for (i = 0; i < 32 * 32; i++) arr[i] = i; - #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + #pragma acc kernels copy(arr) { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { @@ -45,8 +50,8 @@ main (int argc, char* argv[]) arr[i * 32 + j] += x; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32 * 32; i++) assert (arr[i] == i + ((i / 32) ^ (i % 32) * 3)); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c index 80ac99013d6..40baae34562 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared on a loop directive, broadcasting to vector-partitioned mode. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,17 +30,25 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -49,8 +61,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += x * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c index a05ac609123..c8b089c0f59 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared on a loop directive, broadcasting to vector-partitioned mode. Back-to-back worker loops. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,17 +30,25 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -51,7 +63,8 @@ main (int argc, char* argv[]) #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -64,8 +77,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += x * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c index d46bb948626..c1819d27179 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared on a loop directive, broadcasting to vector-partitioned mode. Successive vector loops. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,17 +30,25 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -56,8 +68,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += x * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c index 644c6175863..90955aa8f6e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared on a loop directive, broadcasting to vector-partitioned mode. Addressable worker variable. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -26,17 +30,26 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ + /* { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_j$c_loop_j } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ /* { dg-note {variable 'p' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) @@ -59,8 +72,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += x * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c index 182a12a1b3f..f093cfe630b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared on a loop directive, broadcasting to vector-partitioned mode. Aggregate worker variable. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -32,17 +36,25 @@ main (int argc, char* argv[]) arr[i] = i; #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { #pragma acc loop worker(num:32) private(pt) /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -62,8 +74,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += pt.y * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c index bdfbb59c5c5..906119caf24 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c @@ -1,11 +1,15 @@ /* Test of worker-private variables declared on loop directive, broadcasting to vector-partitioned mode. Array worker variable. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */ + { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -29,18 +33,26 @@ main (int argc, char* argv[]) /* "pt" is treated as "present_or_copy" on the kernels directive because it is an array variable. */ #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition + decides that a data region is needed for 'j', and subsequently requests it + to be made addressable. + { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute } + { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */ { int j; + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 32; i++) { /* But here, it is made private per-worker. */ #pragma acc loop worker(num:32) private(pt) /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ - /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ + /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j } + { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */ for (j = 0; j < 32; j++) { @@ -61,8 +73,8 @@ main (int argc, char* argv[]) arr[i * 1024 + j * 32 + k] += pt[1] * k; } } + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ } - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ for (i = 0; i < 32; i++) for (int j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90 index 09ab3956624..b4f44d69de1 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90 @@ -2,11 +2,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -23,7 +27,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } @@ -31,8 +36,8 @@ program main x = i * 2; arr(i) = arr(i) + x; end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 1, 32 if (arr(i) .ne. i * 3) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90 index bec1069c2a8..f4194c96f90 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,7 +28,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } @@ -37,8 +42,8 @@ program main arr(i * 32 + j) = arr(i * 32 + j) + x; end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 * 32 - 1 if (arr(i) .ne. i + (i / 32) * 2) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90 index 9fde012c19c..da85ac485ad 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,7 +28,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } @@ -37,8 +42,8 @@ program main arr(i * 32 + j) = arr(i * 32 + j) + x; end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 * 32 - 1 if (arr(i) .ne. i + (i / 32) * 2) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90 index 02e09b31cf7..b2bcf11cef0 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -29,7 +33,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) private(pt) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } ! { dg-note {variable 'pt' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } @@ -45,8 +50,8 @@ program main arr(i * 32 + j) = arr(i * 32 + j) + pt%x + pt%y + pt%z + pt%attr(5); end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 * 32 - 1 if (arr(i) .ne. i + (i / 32) * 13) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90 index 5811d0c41b7..eba7263ae56 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90 @@ -2,11 +2,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -23,7 +27,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -46,8 +51,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90 index 81125a24d0d..7d1c6885cee 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90 @@ -2,11 +2,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -23,7 +27,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -41,8 +46,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90 index 824c198799c..d459a6f185e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90 @@ -2,11 +2,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -25,6 +29,8 @@ program main end do !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -36,8 +42,8 @@ program main arr(i * 32 + j) = arr(i * 32 + j) + x end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 * 32 - 1 if (arr(i) .ne. i + ieor(i / 32, mod(i, 32) * 3)) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90 index d25d419a11e..c3a4e2a3406 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,7 +28,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -41,8 +46,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90 index 7a69145b3b6..db6097b6fd5 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,7 +28,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -54,8 +59,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90 index 2c1d5665ce5..a0b6c29d466 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,7 +28,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -49,8 +54,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90 index 4936e569b44..22c0b9b1045 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -27,6 +31,8 @@ program main end do !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -53,8 +59,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90 index 6b2ec1a0047..1b4d5036596 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -29,7 +33,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -53,8 +58,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90 index a90be1d45e8..9ee8a5afc2e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90 @@ -3,11 +3,15 @@ ! { dg-do run } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "-foffload=-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting: +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' ! passed to 'incr' may be unset, and in that case, it will be set to [...]", @@ -24,7 +28,8 @@ program main arr(i) = i end do - !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copy(arr) + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } do i = 0, 31 @@ -48,8 +53,8 @@ program main end do end do end do + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } !$acc end kernels - ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } do i = 0, 32 - 1 do j = 0, 32 -1 -- 2.34.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
* (no subject) 2022-03-04 13:46 ` OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] Thomas Schwinge 2022-03-04 14:57 ` Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784] Thomas Schwinge @ 2022-03-04 19:52 ` Thomas Schwinge 1 sibling, 0 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-04 19:52 UTC (permalink / raw) To: gcc-patches Cc: fortran, Jakub Jelinek, Julian Brown, asolokha, Andrew Pinski, seurer [-- Attachment #1: Type: text/plain, Size: 2576 bytes --] Hi! On 2022-03-04T14:46:25+0100, I wrote: > Pushed to master branch commit 8935589b496f755e08cadf26d8ceddf0dd6e0968 > "OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs > [PR100280, PR104132, PR104133]", see attached. > --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c > [...] > @@ -27,8 +31,12 @@ int main() > (volatile int *) &a; > #define N 123 > int b[N] = { 0 }; > + unsigned long long f1; > + /*TODO See above. */ > + (volatile void *) &f1; Ah, the famous last-minute change just before 'git push'... To work around execution failure with GCN offloading, we're explicitly making 'f1' addressable here -- but I didn't realize that this also affects diagnostics, sorry. Pushed to master branch commit 14dfbb53594e164fe222476523a68039a8bd5252 "Fix 'libgomp.oacc-c-c++-common/kernels-decompose-1.c' expected diagnostics", see attached. Grüße Thomas > > #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ > + /* { dg-note {variable 'g2\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ > { > [...] > + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ > + f1 = 1; > + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ > +#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */ > + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */ > + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */ > + for (c = 20; c > 0; --c) > + f1 *= c; > + > + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ > + if (c != 234) > + __builtin_abort (); > + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ > + } > } > [...] > + assert (f1 == 2432902008176640000ULL); > > return 0; > } ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Fix-libgomp.oacc-c-c-common-kernels-decompose-1.c-ex.patch --] [-- Type: text/x-diff, Size: 1827 bytes --] From 14dfbb53594e164fe222476523a68039a8bd5252 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 4 Mar 2022 20:34:40 +0100 Subject: [PATCH] Fix 'libgomp.oacc-c-c++-common/kernels-decompose-1.c' expected diagnostics Fix-up for recent commit 8935589b496f755e08cadf26d8ceddf0dd6e0968 "OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133]": adjust for a GCN offloading workaround added just before commit: '(volatile void *) &f1;'. PR testsuite/104791 libgomp/ * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Fix expected diagnostics. --- .../testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c | 2 ++ 1 file changed, 2 insertions(+) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index 049b3a44b03..985a547d381 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -37,6 +37,8 @@ int main() #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ /* { dg-note {variable 'g2\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'f1\.1' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'f1\.2' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ int c = 234; -- 2.25.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280] 2022-01-13 9:54 ` OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] Thomas Schwinge 2022-03-01 16:46 ` Thomas Schwinge @ 2022-03-04 13:37 ` Thomas Schwinge 2022-03-04 13:41 ` OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280] Thomas Schwinge 2 siblings, 0 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-04 13:37 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek, Julian Brown, asolokha [-- Attachment #1: Type: text/plain, Size: 2615 bytes --] Hi! On 2022-01-13T10:54:16+0100, I wrote: > On 2019-05-08T14:51:57+0100, Julian Brown <julian@codesourcery.com> wrote: >> - The "addressable" bit is set during the kernels conversion pass for >> variables that have "create" (alloc) clauses created for them in the >> synthesised outer data region (instead of in the front-end, etc., >> where it can't be done accurately). Such variables actually have >> their address taken during transformations made in a later pass >> (omp-low, I think), but there's a phase-ordering problem that means >> the flag should be set earlier. > > The actual issue is a bit different, but yes, there is a problem. > The related ICE has also been reported as <https://gcc.gnu.org/PR100280> > "ICE in lower_omp_target, at omp-low.c:12287". (And I'm confused why we > didn't run into that with the OpenACC 'kernels' decomposition > originally.) I've pushed to master branch > commit 9b32c1669aad5459dd053424f9967011348add83 > "OpenACC 'kernels' decomposition: Mark variables used in synthesized data > clauses as addressable [PR100280]", see attached. > --- a/gcc/omp-oacc-kernels-decompose.cc > +++ b/gcc/omp-oacc-kernels-decompose.cc > @@ -793,7 +793,8 @@ make_data_region_try_statement (location_t loc, gimple *body) > > /* If INNER_BIND_VARS holds variables, build an OpenACC data region with > location LOC containing BODY and having 'create (var)' clauses for each > - variable. If INNER_CLEANUP is present, add a try-finally statement with > + variable (as a side effect, such variables also get TREE_ADDRESSABLE set). > + If INNER_CLEANUP is present, add a try-finally statement with > this cleanup code in the finally block. Return the new data region, or > the original BODY if no data region was needed. */ > > @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, > inner_data_clauses = new_clause; > > prev_mapped_var = v; > + > + /* See <https://gcc.gnu.org/PR100280>. */ > + TREE_ADDRESSABLE (v) = 1; > } > } Pushed to master branch commit e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc "Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280]", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Add-diagnostic-note-OpenACC-kernels-decomposition-va.patch --] [-- Type: text/x-diff, Size: 16806 bytes --] From e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Tue, 15 Feb 2022 16:54:30 +0100 Subject: [PATCH] Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280] Follow-up to commit 9b32c1669aad5459dd053424f9967011348add83 "OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280]". PR middle-end/100280 gcc/ * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region): Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable". gcc/testsuite/ * c-c++-common/goacc/classify-kernels-unparallelized.c: Add '--param=openacc-privatization=noisy'. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/kernels-decompose-2.c: Adjust. * c-c++-common/goacc/kernels-decompose-pr100280-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. --- gcc/omp-oacc-kernels-decompose.cc | 24 ++++++++++++++++++- .../goacc/classify-kernels-unparallelized.c | 7 ++++++ .../c-c++-common/goacc/classify-kernels.c | 7 ++++++ .../c-c++-common/goacc/kernels-decompose-2.c | 2 ++ .../goacc/kernels-decompose-pr100280-1.c | 1 + .../goacc/kernels-decompose-pr104061-1-2.c | 1 + .../goacc/kernels-decompose-pr104061-1-3.c | 1 + .../goacc/kernels-decompose-pr104061-1-4.c | 1 + .../goacc/kernels-decompose-pr104132-1.c | 1 + .../goacc/kernels-decompose-pr104133-1.c | 1 + .../libgomp.oacc-c-c++-common/f-asyncwait-1.c | 3 +++ .../kernels-decompose-1.c | 1 + 12 files changed, 49 insertions(+), 1 deletion(-) diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index 98eafdbe3a1..5093386f718 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -845,7 +845,29 @@ maybe_build_inner_data_region (location_t loc, gimple *body, prev_mapped_var = v; /* See <https://gcc.gnu.org/PR100280>. */ - TREE_ADDRESSABLE (v) = 1; + if (!TREE_ADDRESSABLE (v)) + { + TREE_ADDRESSABLE (v) = 1; + + if (dump_enabled_p ()) + { + const dump_user_location_t d_u_loc + = dump_user_location_t::from_location_t (loc); + /* PR100695 "Format decoder, quoting in 'dump_printf' etc." */ +#if __GNUC__ >= 10 +# pragma GCC diagnostic push +# pragma GCC diagnostic ignored "-Wformat" +#endif + dump_printf_loc (MSG_NOTE, d_u_loc, + "OpenACC %<kernels%> decomposition:" + " variable %<%T%> declared in block" + " made addressable\n", + v); +#if __GNUC__ >= 10 +# pragma GCC diagnostic pop +#endif + } + } } } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index 4ee8e9d5f39..2496462f777 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -9,6 +9,10 @@ { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccloops" } */ +/* { dg-additional-options "--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ + /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -25,6 +29,9 @@ extern unsigned int f (unsigned int); void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute1 } */ /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (unsigned int i = 0; i < N; i++) diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 74acca8b4a6..3ba0411a57a 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -9,6 +9,10 @@ { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccloops" } */ +/* { dg-additional-options "--param=openacc-privatization=noisy" } + Prune a few: uninteresting: + { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ + /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -21,6 +25,9 @@ extern unsigned int *__restrict c; void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute1 } */ /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (unsigned int i = 0; i < N; i++) diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c index 64ce8943a02..5fbd10221a2 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -66,12 +66,14 @@ main () #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ { int i; } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c index b497af298bc..91e589777b0 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c @@ -10,6 +10,7 @@ void foo (void) /* { dg-line l_f_1 } */ { #pragma acc kernels /* { dg-line l_k_1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_k_1 } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_k_1 } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_k_1 } */ /* { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {TODO location} { xfail *-*-* } l_f_1 } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c index ea58ea4c161..8ca02cb2119 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c @@ -16,6 +16,7 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_compute1 } */ { /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c index ee3fcc7c39a..05a196d8f36 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c @@ -23,6 +23,7 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c index c140e61e3b5..07cb592bcf9 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c @@ -21,6 +21,7 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c index 4fbfdd81e15..5b0fe42efbf 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c @@ -18,6 +18,7 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ { int k; diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c index 72dde346dbf..4536b5c179e 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c @@ -18,6 +18,7 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.1' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c index ef7735b2ef4..d360aad4736 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c @@ -260,6 +260,7 @@ main (void) #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ @@ -268,6 +269,7 @@ main (void) b[i] = (a[i] * a[i] * a[i]) / a[i]; #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ @@ -311,6 +313,7 @@ main (void) #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index 57e75f6d399..b99497ea6c5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -32,6 +32,7 @@ int main() { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ int c = 234; + /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ #pragma acc loop independent gang /* { dg-line l_loop_i[incr c_loop_i] } */ -- 2.34.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
* OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280] 2022-01-13 9:54 ` OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] Thomas Schwinge 2022-03-01 16:46 ` Thomas Schwinge 2022-03-04 13:37 ` Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280] Thomas Schwinge @ 2022-03-04 13:41 ` Thomas Schwinge 2 siblings, 0 replies; 10+ messages in thread From: Thomas Schwinge @ 2022-03-04 13:41 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek, Julian Brown, asolokha [-- Attachment #1: Type: text/plain, Size: 2589 bytes --] Hi! On 2022-01-13T10:54:16+0100, I wrote: > On 2019-05-08T14:51:57+0100, Julian Brown <julian@codesourcery.com> wrote: >> - The "addressable" bit is set during the kernels conversion pass for >> variables that have "create" (alloc) clauses created for them in the >> synthesised outer data region (instead of in the front-end, etc., >> where it can't be done accurately). Such variables actually have >> their address taken during transformations made in a later pass >> (omp-low, I think), but there's a phase-ordering problem that means >> the flag should be set earlier. > > The actual issue is a bit different, but yes, there is a problem. > The related ICE has also been reported as <https://gcc.gnu.org/PR100280> > "ICE in lower_omp_target, at omp-low.c:12287". (And I'm confused why we > didn't run into that with the OpenACC 'kernels' decomposition > originally.) I've pushed to master branch > commit 9b32c1669aad5459dd053424f9967011348add83 > "OpenACC 'kernels' decomposition: Mark variables used in synthesized data > clauses as addressable [PR100280]", see attached. > --- a/gcc/omp-oacc-kernels-decompose.cc > +++ b/gcc/omp-oacc-kernels-decompose.cc > @@ -793,7 +793,8 @@ make_data_region_try_statement (location_t loc, gimple *body) > > /* If INNER_BIND_VARS holds variables, build an OpenACC data region with > location LOC containing BODY and having 'create (var)' clauses for each > - variable. If INNER_CLEANUP is present, add a try-finally statement with > + variable (as a side effect, such variables also get TREE_ADDRESSABLE set). > + If INNER_CLEANUP is present, add a try-finally statement with > this cleanup code in the finally block. Return the new data region, or > the original BODY if no data region was needed. */ > > @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, > inner_data_clauses = new_clause; > > prev_mapped_var = v; > + > + /* See <https://gcc.gnu.org/PR100280>. */ > + TREE_ADDRESSABLE (v) = 1; > } > } Pushed to master branch commit de6e81ea961219d0726db67776d11ce75a4cae1b "OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280]", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-OpenACC-kernels-decomposition-Move-TREE_ADDRESSABLE-.patch --] [-- Type: text/x-diff, Size: 21643 bytes --] From de6e81ea961219d0726db67776d11ce75a4cae1b Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Tue, 15 Feb 2022 23:03:49 +0100 Subject: [PATCH] OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280] ... in preparation for later changes. No functional change. Follow-up to commit 9b32c1669aad5459dd053424f9967011348add83 "OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280]". PR middle-end/100280 gcc/ * tree.h (OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE): New. * tree-core.h: Document it. * omp-low.cc (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Handle 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'. * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region): Set 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' instead of 'TREE_ADDRESSABLE'. gcc/testsuite/ * c-c++-common/goacc/classify-kernels-unparallelized.c: Adjust. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/kernels-decompose-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr100280-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. --- gcc/omp-low.cc | 31 +++++++++++++++++++ gcc/omp-oacc-kernels-decompose.cc | 5 +-- .../goacc/classify-kernels-unparallelized.c | 3 +- .../c-c++-common/goacc/classify-kernels.c | 3 +- .../c-c++-common/goacc/kernels-decompose-2.c | 6 ++-- .../goacc/kernels-decompose-pr100280-1.c | 3 +- .../goacc/kernels-decompose-pr104061-1-2.c | 3 +- .../goacc/kernels-decompose-pr104061-1-3.c | 3 +- .../goacc/kernels-decompose-pr104061-1-4.c | 3 +- .../goacc/kernels-decompose-pr104132-1.c | 3 +- .../goacc/kernels-decompose-pr104133-1.c | 3 +- gcc/tree-core.h | 3 ++ gcc/tree.h | 5 +++ .../libgomp.oacc-c-c++-common/f-asyncwait-1.c | 9 ++++-- .../kernels-decompose-1.c | 3 +- 15 files changed, 70 insertions(+), 16 deletions(-) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 2294456b27d..6654bfd426e 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1495,6 +1495,37 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); decl = OMP_CLAUSE_DECL (c); + /* If requested, make 'decl' addressable. */ + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c)) + { + gcc_checking_assert (DECL_P (decl)); + + gcc_checking_assert (!TREE_ADDRESSABLE (decl)); + TREE_ADDRESSABLE (decl) = 1; + + if (dump_enabled_p ()) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + const dump_user_location_t d_u_loc + = dump_user_location_t::from_location_t (loc); + /* PR100695 "Format decoder, quoting in 'dump_printf' etc." */ +#if __GNUC__ >= 10 +# pragma GCC diagnostic push +# pragma GCC diagnostic ignored "-Wformat" +#endif + dump_printf_loc (MSG_NOTE, d_u_loc, + "variable %<%T%>" + " made addressable\n", + decl); +#if __GNUC__ >= 10 +# pragma GCC diagnostic pop +#endif + } + + /* Done. */ + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c) = 0; + } /* Global variables with "omp declare target" attribute don't need to be copied, the receiver side will use them directly. However, global variables with "omp declare target link" diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index 5093386f718..ecbd3071e5d 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -847,7 +847,8 @@ maybe_build_inner_data_region (location_t loc, gimple *body, /* See <https://gcc.gnu.org/PR100280>. */ if (!TREE_ADDRESSABLE (v)) { - TREE_ADDRESSABLE (v) = 1; + /* Request that OMP lowering make 'v' addressable. */ + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1; if (dump_enabled_p ()) { @@ -861,7 +862,7 @@ maybe_build_inner_data_region (location_t loc, gimple *body, dump_printf_loc (MSG_NOTE, d_u_loc, "OpenACC %<kernels%> decomposition:" " variable %<%T%> declared in block" - " made addressable\n", + " requested to be made addressable\n", v); #if __GNUC__ >= 10 # pragma GCC diagnostic pop diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index 2496462f777..61871d118a9 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -29,7 +29,8 @@ extern unsigned int f (unsigned int); void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 3ba0411a57a..1473337986f 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -25,7 +25,8 @@ extern unsigned int *__restrict c; void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c index 5fbd10221a2..bf158311dae 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -66,14 +66,16 @@ main () #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ { int i; } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c index 91e589777b0..1c1e22c00ac 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c @@ -10,7 +10,8 @@ void foo (void) /* { dg-line l_f_1 } */ { #pragma acc kernels /* { dg-line l_k_1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_k_1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_k_1 } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_k_1 } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_k_1 } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_k_1 } */ /* { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {TODO location} { xfail *-*-* } l_f_1 } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c index 8ca02cb2119..336cf2ad425 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c @@ -16,7 +16,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_compute1 } */ { /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c index 05a196d8f36..f41dda86122 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c @@ -23,7 +23,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c index 07cb592bcf9..cde95a7b7ac 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c @@ -21,7 +21,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c index 5b0fe42efbf..4f38a83bb19 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c @@ -18,7 +18,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'k' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ { int k; diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c index 4536b5c179e..0499665777d 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c @@ -18,7 +18,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'k' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.1' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 01a1ce499da..4530bd8c2c7 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -1155,6 +1155,9 @@ struct GTY(()) tree_base { PREDICT_EXPR_OUTCOME in PREDICT_EXPR + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE in + OMP_CLAUSE + static_flag: TREE_STATIC in diff --git a/gcc/tree.h b/gcc/tree.h index 36ceed57064..da6f3b38ba5 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1699,6 +1699,11 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag) +/* Flag that 'OMP_CLAUSE_DECL (NODE)' is to be made addressable during OMP + lowering. */ +#define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag) + /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' clause. */ #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c index d360aad4736..9f7b8b4af2f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c @@ -260,7 +260,8 @@ main (void) #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ @@ -269,7 +270,8 @@ main (void) b[i] = (a[i] * a[i] * a[i]) / a[i]; #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ @@ -313,7 +315,8 @@ main (void) #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index b99497ea6c5..85c39871f94 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -32,7 +32,8 @@ int main() { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ int c = 234; - /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'c' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ #pragma acc loop independent gang /* { dg-line l_loop_i[incr c_loop_i] } */ -- 2.34.1 ^ permalink raw reply [flat|nested] 10+ messages in thread
end of thread, other threads:[~2022-03-04 19:52 UTC | newest]
Thread overview: 10+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <20190508145157.08beb4df@squid.athome>
2022-01-13 9:54 ` OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] Thomas Schwinge
2022-03-01 16:46 ` Thomas Schwinge
2022-03-02 13:30 ` Jakub Jelinek
2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104132-1.c' [PR104132] Thomas Schwinge
2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104133-1.c' [PR104133] Thomas Schwinge
2022-03-04 13:46 ` OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] Thomas Schwinge
2022-03-04 14:57 ` Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784] Thomas Schwinge
2022-03-04 19:52 ` Thomas Schwinge
2022-03-04 13:37 ` Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280] Thomas Schwinge
2022-03-04 13:41 ` OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280] Thomas Schwinge
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).