From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id B4FA73857C50; Fri, 4 Mar 2022 13:26:05 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org B4FA73857C50 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-7481] OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/master X-Git-Oldrev: de6e81ea961219d0726db67776d11ce75a4cae1b X-Git-Newrev: 8935589b496f755e08cadf26d8ceddf0dd6e0968 Message-Id: <20220304132605.B4FA73857C50@sourceware.org> Date: Fri, 4 Mar 2022 13:26:05 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Fri, 04 Mar 2022 13:26:05 -0000 https://gcc.gnu.org/g:8935589b496f755e08cadf26d8ceddf0dd6e0968 commit r12-7481-g8935589b496f755e08cadf26d8ceddf0dd6e0968 Author: Thomas Schwinge Date: Tue Feb 15 23:31:34 2022 +0100 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) 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. Diff: --- 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 taskreg_contexts; static vec 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 (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 (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 +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; }