From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from us-smtp-delivery-124.mimecast.com (us-smtp-delivery-124.mimecast.com [170.10.129.124]) by sourceware.org (Postfix) with ESMTPS id A41FA3858D39 for ; Thu, 3 Mar 2022 08:24:10 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org A41FA3858D39 Received: from mimecast-mx01.redhat.com (mimecast-mx01.redhat.com [209.132.183.4]) by relay.mimecast.com with ESMTP with STARTTLS (version=TLSv1.2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id us-mta-665-gx9BRgKuP4OS7FtHeCza_w-1; Thu, 03 Mar 2022 03:24:06 -0500 X-MC-Unique: gx9BRgKuP4OS7FtHeCza_w-1 Received: from smtp.corp.redhat.com (int-mx07.intmail.prod.int.phx2.redhat.com [10.5.11.22]) (using TLSv1.2 with cipher AECDH-AES256-SHA (256/256 bits)) (No client certificate requested) by mimecast-mx01.redhat.com (Postfix) with ESMTPS id 2E207801AAD; Thu, 3 Mar 2022 08:24:05 +0000 (UTC) Received: from tucnak.zalov.cz (unknown [10.39.192.81]) by smtp.corp.redhat.com (Postfix) with ESMTPS id BFE4A1059589; Thu, 3 Mar 2022 08:24:04 +0000 (UTC) Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.16.1/8.16.1) with ESMTPS id 2238O16Q3001061 (version=TLSv1.3 cipher=TLS_AES_256_GCM_SHA384 bits=256 verify=NOT); Thu, 3 Mar 2022 09:24:02 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.16.1/8.16.1/Submit) id 2238O1HP3001060; Thu, 3 Mar 2022 09:24:01 +0100 Date: Thu, 3 Mar 2022 09:24:01 +0100 From: Jakub Jelinek To: gcc-patches@gcc.gnu.org Cc: Tobias Burnus Subject: [committed] openmp: Disable SSA form during gimplification on OMP_SIMD clauses and body [PR104757] Message-ID: Reply-To: Jakub Jelinek MIME-Version: 1.0 X-Scanned-By: MIMEDefang 2.84 on 10.5.11.22 X-Mimecast-Spam-Score: 0 X-Mimecast-Originator: redhat.com Content-Type: text/plain; charset=us-ascii Content-Disposition: inline X-Spam-Status: No, score=-5.3 required=5.0 tests=BAYES_00, DKIMWL_WL_HIGH, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, RCVD_IN_DNSWL_LOW, RCVD_IN_MSPIKE_H5, RCVD_IN_MSPIKE_WL, SPF_HELO_NONE, SPF_NONE, 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: , X-List-Received-Date: Thu, 03 Mar 2022 08:24:12 -0000 Hi! When offloading to nvptx is enabled, scan_omp_simd duplicates the simd region including its clauses and body using inliner's copy_gimple_seq_and_replace_locals. That works nicely for decls, remaps only those that are seen in the nested bind expr vars (i.e. local variables) and doesn't remap other vars. But for SSA_NAMEs it remaps them always, doesn't know if their def stmt is outside of the simd (then it better shouldn't be remapped) or inside of it (then it should) and without cfg/dominators that is pretty hard to figure out (well, we could walk the region twice, once note SSA_NAMEs defined by each stmt seen there and once do the remapping of only those visited SSA_NAMEs). This patch uses a simpler way, disables temporarily into_ssa for the clauses and body of each simd region; we already disable into_ssa e.g. in parallel/target/task etc. regions through push_gimplify_context () but for simd we don't push any gimplification context and appart from into_ssa I think we don't need it. Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk. 2022-03-03 Jakub Jelinek PR middle-end/104757 * gimplify.cc (gimplify_omp_loop): Call gimplify_expr rather than gimplify_omp_for. (gimplify_expr) : Temporarily disable gimplify_ctxp->into_ssa around call to gimplify_omp_for. * gfortran.dg/gomp/pr104757.f90: New test. * gcc.dg/gomp/pr104757.c: New test. --- gcc/gimplify.cc.jj 2022-02-14 13:14:39.058128337 +0100 +++ gcc/gimplify.cc 2022-03-02 19:31:46.675210108 +0100 @@ -13725,7 +13725,7 @@ gimplify_omp_loop (tree *expr_p, gimple_ *pc = NULL_TREE; *expr_p = t; } - return gimplify_omp_for (expr_p, pre_p); + return gimplify_expr (expr_p, pre_p, NULL, is_gimple_stmt, fb_none); } @@ -15479,8 +15479,19 @@ gimplify_expr (tree *expr_p, gimple_seq ret = GS_ALL_DONE; break; - case OMP_FOR: case OMP_SIMD: + { + /* Temporarily disable into_ssa, as scan_omp_simd + which calls copy_gimple_seq_and_replace_locals can't deal + with SSA_NAMEs defined outside of the body properly. */ + bool saved_into_ssa = gimplify_ctxp->into_ssa; + gimplify_ctxp->into_ssa = false; + ret = gimplify_omp_for (expr_p, pre_p); + gimplify_ctxp->into_ssa = saved_into_ssa; + break; + } + + case OMP_FOR: case OMP_DISTRIBUTE: case OMP_TASKLOOP: case OACC_LOOP: --- gcc/testsuite/gfortran.dg/gomp/pr104757.f90.jj 2022-03-02 18:35:00.606549074 +0100 +++ gcc/testsuite/gfortran.dg/gomp/pr104757.f90 2022-03-02 18:34:42.028804663 +0100 @@ -0,0 +1,19 @@ +! PR middle-end/104757 +! { dg-do compile } +! { dg-options "-O -fopenmp" } + +module pr104757 + implicit none (external, type) + integer :: ll + !$omp declare target (ll) +contains + subroutine foo (i1) + !$omp declare target (foo) + logical :: i1 + integer :: i + !$omp distribute simd if(i1) + do i = 1, 64 + ll = ll + 1 + end do + end +end module --- gcc/testsuite/gcc.dg/gomp/pr104757.c.jj 2022-03-02 19:33:52.637485113 +0100 +++ gcc/testsuite/gcc.dg/gomp/pr104757.c 2022-03-02 19:34:09.365256034 +0100 @@ -0,0 +1,14 @@ +/* PR middle-end/104757 */ +/* { dg-do compile } */ +/* { dg-options "-O2 -fopenmp" } */ + +#pragma omp declare target +void +foo (int x, int y, int *z) +{ + int j = 0; + #pragma omp simd linear(j:x + y) + for (int i = 0; i < 64; i++) + j += x + y; +} +#pragma omp end declare target Jakub