From patchwork Fri Mar 4 13:46:25 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 51574 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id C0DD73857C74 for ; Fri, 4 Mar 2022 13:47:02 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id D13CD3858D39; Fri, 4 Mar 2022 13:46:36 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org D13CD3858D39 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.90,155,1643702400"; d="scan'208,223";a="75299564" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 04 Mar 2022 05:46:37 -0800 IronPort-SDR: iyCILlZqaoENojKMULkpe9V2ph9EAILUgbgXtH6WKw8pn5VekAckTnQFkVykX0/hfXt1yT+/+X GV6kEvEceMd4vXxMWJzEJhHrdUyVNPyHOGRETRDieaOA28WTmqnZ9+HxqKcDYtV4DuJu7olXVN BI472elirnbAQXy1CizA61xD10ZzVDtklQtGOkJDTjbR+V2SXfXXog3FYcP5BGM8CkpghmLlDE 9YqZ6LaPQWPa+MxGZNGgF5ZBLZw/SMCjDQcnwInWn7w3LMMwahMgzhbMyyXl5ZbqjkQ6/AHG0+ dhI= From: Thomas Schwinge To: Subject: OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] In-Reply-To: <87zgm9mxib.fsf@euler.schwinge.homeip.net> References: <20190508145157.08beb4df@squid.athome> <87iluovu47.fsf@euler.schwinge.homeip.net> <87zgm9mxib.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Fri, 4 Mar 2022 14:46:25 +0100 Message-ID: <877d99omoe.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , Julian Brown , fortran@gcc.gnu.org, asolokha@gmx.com Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" 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 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 >> "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 . */ >> + 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 . > 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 From 8935589b496f755e08cadf26d8ceddf0dd6e0968 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge 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) 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 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; } -- 2.34.1