From patchwork Fri Jun 7 14:01:29 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1111853 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-502568-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="Hlil1Wvy"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 45L4320xnwz9s7h for ; Sat, 8 Jun 2019 00:02:01 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=x+G1xg5zFAvWtEe1 nkUQvFab5/sYicG8AjJ3FaAPP3lKzhka9+36L0rIzB2KjDikPv7c4mdXLofSPK9d on/iI5DSUE6CnaL/RK/FuV++33KEDArNwwwQeviBpjGqO6hkRBJxV+NJKVjeSVLw yIbDdwGMIKXjSZ1RcvpQ0jlvYSc= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=Bjp8rdpDsL2Owi9VCPI9A0 4geTw=; b=Hlil1WvyGY8Jbr4KhGaWLDyyb4jGVuHTfbM/TaukJEwNCZy30FARbG Rx/wD14mrRX1mdWZ9uYRBN/5tvwJAcTlH2WlSAosnDR5Djo8I4SEA8zEAO9e+zlC v27jqUACWpfZ2JW4ms6Lq8vAn5WbuJS1vqIC7rkOgkXXC2/oJeK+8= Received: (qmail 20308 invoked by alias); 7 Jun 2019 14:01:51 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 20293 invoked by uid 89); 7 Jun 2019 14:01:51 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: =?iso-8859-1?q?No=2C_score=3D-7=2E7_required=3D5=2E0?= =?iso-8859-1?q?_tests=3DAWL=2CBAYES_20=2CGIT_PATCH_2=2CGIT_PATCH_3?= =?iso-8859-1?q?=2CKAM_SHORT=2CRCVD_IN_DNSWL_NONE_autolearn=3Dham_v?= =?iso-8859-1?q?ersion=3D3=2E3=2E1_spammy=3Dgr=2C_gr=C3=BC=C3=2C_om?= =?iso-8859-1?q?p_context=2C_nevertheless?= X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 07 Jun 2019 14:01:48 +0000 Received: from svr-orw-mbx-06.mgc.mentorg.com ([147.34.90.206]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1hZFRA-0001LI-Ft from Thomas_Schwinge@mentor.com ; Fri, 07 Jun 2019 07:01:44 -0700 Received: from SVR-ORW-MBX-09.mgc.mentorg.com (147.34.90.209) by SVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Fri, 7 Jun 2019 07:01:42 -0700 Received: from tftp-cs (147.34.91.1) by SVR-ORW-MBX-09.mgc.mentorg.com (147.34.90.209) with Microsoft SMTP Server id 15.0.1320.4 via Frontend Transport; Fri, 7 Jun 2019 07:01:42 -0700 Received: by tftp-cs (Postfix, from userid 49978) id B07C9C230F; Fri, 7 Jun 2019 07:01:41 -0700 (PDT) From: Thomas Schwinge To: , Jakub Jelinek CC: Subject: [PR90742] OpenACC/OpenMP target offloading: Fortran 'allocatable' scalars in 'firstprivate' clauses In-Reply-To: <1f88e441-d3da-5b59-4278-058ff1368a73@codesourcery.com> References: <1f88e441-d3da-5b59-4278-058ff1368a73@codesourcery.com> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Fri, 7 Jun 2019 16:01:29 +0200 Message-ID: <87r285h4ue.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! As I had mentioned in the PR... On Tue, 7 Aug 2018 14:55:07 -0700, Cesar Philippidis wrote: > This patch ... would be one component for fixing "OpenACC/OpenMP target offloading: Fortran 'allocatable' scalars in 'firstprivate' clauses". (Also, as mentioned there, such changes have been submitted already, a few times, muddled into other changes. So, thanks, that this also got submitted separately, to address just this one issue.) > updates the way that lower_omp_target uses firstprivate > pointers in OpenACC offloaded regions. On host side, when preparing > firstprivate data mapping for pointer type objects, not to be confused > with GOMP_MAP_FIRSTPRIVATE_POINTER, the compiler passes passes the > address of the value being pointed to and not the address of the pointer > itself to the runtime. Correspondingly, on the device side, the compiler > generates to code to dereference the remapped pointer once to copy the > data to a local buffer. > > While this behavior looks like it would break things, it will not affect > C or C++ data mappings, because those languages transfer pointers via > GOMP_MAP_FIRSTPRIVATE_POINTER. Not with current GCC sources, as I should eventually find out, which are still missing another patch or two, or three, or more. > In addition, this will not cause > problems with array types, because the default remapping rules for > OpenACC is to transfer them in via copy. Besides it really doesn't > make sense to allow arrays to be transferred in via firstprivate > because that would use up a lot of memory on the accelerator. (Huh, but the latter ought to be supported nevertheless, as far as I understand? Anyway, that'll be for later.) > Is this OK for trunk? I bootstrapped and regtested it for x86_64 with > nvptx offloading. The patch, as proposed, does introduce regressions. > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -7643,15 +7643,21 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) > { > gcc_assert (is_gimple_omp_oacc (ctx->stmt)); > - if (omp_is_reference (new_var) > - && TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE) > + if (omp_is_reference (new_var)) > { > /* Create a local object to hold the instance > value. */ > - tree type = TREE_TYPE (TREE_TYPE (new_var)); > + tree type = TREE_TYPE (new_var); > + /* Pointer types are mapped onto the device via a > + single level of indirection. */ > + if (TREE_CODE (type) != POINTER_TYPE) > + type = TREE_TYPE (type); > const char *id = IDENTIFIER_POINTER (DECL_NAME (new_var)); > tree inst = create_tmp_var (type, id); > - gimplify_assign (inst, fold_indirect_ref (x), &fplist); > + if (TREE_CODE (TREE_TYPE (new_var)) == POINTER_TYPE) > + gimplify_assign (inst, fold_indirect_ref (x), &fplist); > + else > + gimplify_assign (inst, fold_indirect_ref (x), &fplist); > x = build_fold_addr_expr (inst); > } > gimplify_assign (new_var, x, &fplist); (It seems strange to have the same code in both branches of the 'if' statement?) > @@ -7879,7 +7885,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) > { > gcc_assert (is_gimple_omp_oacc (ctx->stmt)); > - if (!omp_is_reference (var)) > + /* Handle Fortran allocatable scalars. */ > + if (!omp_is_reference (var) > + && TREE_CODE (TREE_TYPE (var)) != POINTER_TYPE) > { > if (is_gimple_reg (var) > && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)) | TREE_NO_WARNING (var) = 1; | var = build_fold_addr_expr (var); | } | else | talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar))); | gimplify_assign (x, var, &ilist); | } That's what's causing regressions, for example for 'firstprivate' clauses even in non-offloading situation ('if(0)' clause, for example): Program received signal SIGSEGV, Segmentation fault. 0x0000000000402f8a in main._omp_fn.1 () at source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c:59 59 b[ii] = a[ii] + 1; (gdb) list 10, 10 10 float *a, *b, *d_a, *d_b, exp, exp2; (gdb) list 16, 17 16 a = (float *) malloc (N * sizeof (float)); 17 b = (float *) malloc (N * sizeof (float)); (gdb) list 52, 63 52 #pragma acc parallel if(0) 53 { 54 int ii; 55 56 for (ii = 0; ii < N; ii++) 57 { 58 if (acc_on_device (acc_device_host)) 59 b[ii] = a[ii] + 1; 60 else 61 b[ii] = a[ii]; 62 } 63 } So we got here implicit 'firstprivate(a, b)' (which in this scenario means no-op, given that the host pointer values are just passed through). (On x86_64) these used to have eight bytes alignment, now they have four bytes. But worse, the code on the "sending" side is changed as follows ('omplower' dump): - b.57 = b; - .omp_data_arr.54.b = &b.57; + .omp_data_arr.54.b = b; - a.58 = a; - .omp_data_arr.54.a = &a.58; + .omp_data_arr.54.a = a; #pragma omp target oacc_parallel if(0) firstprivate(b) firstprivate(a) [child fn: main._omp_fn.1 (.omp_data_arr.54, .omp_data_sizes.55, .omp_data_kinds.56)] ..., but the "receiving" side stays the same, so we got a mismatch. If something like that, then the 'POINTER_TYPE' conditional should probably be inside the '!omp_is_reference' conditional, just guarding the 'build_fold_addr_expr'? Anyway, I had a look at this now, and seem to have gotten it work. I will admit, though, that I'm somewhat lost especially with all the 'omp_is_reference' usage ("should privatize what this DECL points to rather than the DECL itself"). Using that on 'OMP_CLAUSE_DECL ([...])' (the common case) makes sense given that's in context of the originating source language, but what exactly does it mean when 'omp_is_reference' is used on 'new_var = lookup_decl (var, ctx)', or on 'var = lookup_decl_in_outer_ctx (ovar, ctx)', where the things looked up by these (that is, stored in 'ctx->cb.decl_map') are "arbitrary"/"synthetic" items? (Jakub?) Or is it actually improper to use 'omp_is_reference' on these, but it just happens to do the expected things in the (several) existing cases? Anyway, for an 'integer, allocatable :: a' that is used 'firstprivate' inside an OpenACC offloading region, we now get the following 'omplower' changes: [...] integer(kind=4) * a; [...] a = __builtin_malloc (4); [...] - a.16 = a; - .omp_data_arr.13.a = &a.16; + .omp_data_arr.13.a = a; #pragma omp target oacc_parallel map(tofrom:b [len: 400]) firstprivate(a) [child fn: MAIN__._omp_fn.0 (.omp_data_arr.13, .omp_data_sizes.14, .omp_data_kinds.15)] { .omp_data_i = (const struct .omp_data_t.10 & restrict) &.omp_data_arr.13; - D.3981 = .omp_data_i->a; - a = *D.3981; + a = .omp_data_i->a; [...] ..., and that seems to work fine. (..., and no testsuite regressions.) (I have not yet looked into the related OpenMP changes required.) Jakub, is the following about right, do you have any comments? If approving this patch, please respond with "Reviewed-by: NAME " so that your effort will be recorded in the commit log, see . The code changes seem very ad-hoc, but that's the common impression that I got from looking at/working though a lot of all that OMP code... :-( Grüße Thomas --- gcc/omp-low.c +++ gcc/omp-low.c @@ -9685,7 +9685,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { gcc_assert (is_gimple_omp_oacc (ctx->stmt)); if (omp_is_reference (new_var) - && TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE) + && TREE_CODE (TREE_TYPE (new_var)) == POINTER_TYPE) + { + /* Special handling for Fortran 'allocatable' scalars: + avoid indirection. */ + x = build_receiver_ref (var, false, ctx); + } + else if (omp_is_reference (new_var)) { /* Create a local object to hold the instance value. */ @@ -9920,7 +9926,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) { gcc_assert (is_gimple_omp_oacc (ctx->stmt)); - if (!omp_is_reference (var)) + if (omp_is_reference (lookup_decl (ovar, ctx)) + && TREE_CODE (TREE_TYPE (ovar)) == POINTER_TYPE) + { + /* Special handling for Fortran 'allocatable' scalars: + avoid indirection. */ + } + else if (!omp_is_reference (var)) { if (is_gimple_reg (var) && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c))