From patchwork Mon Mar 14 23:17:54 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 597273 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 3qPDFs1RnVz9sRB for ; Tue, 15 Mar 2016 10:18:40 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=OOvU3yCY; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=mHL0F5UmsfTf/j9zR QsWU/YfqCgRn9ygbzxcf6dSwLoQZ102Y3/jpdybHTzY6DCMoTXRhyseI4hu401hn FTjXu5BAK9kYucsGuXOSzPr43QwmCdgnRE6dyufc/IoRRWrX0zBqKAK1LRYc98Wl a2sqQGzGmvS5NE80fl6yTSOU3E= 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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=/kT1ULySLQABqw3xoRqIJMx 4cm4=; b=OOvU3yCYZzK8N2Qql9srHQ2ILJNLkwFZW8tSopOsPheHJHZHe/kz5Lf wmSck3mGPOLFmCXAC2qHShrUR1NMD27MzN5CfbGKWr9mIhQvYrWfJRGh13DtWf5q is94zlqOCGLFBang+mO1x2HZEsLrEtyEIJ45uogQ34WTgcqaLy/U= Received: (qmail 13035 invoked by alias); 14 Mar 2016 23:18:28 -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 13018 invoked by uid 89); 14 Mar 2016 23:18:27 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00, RP_MATCHES_RCVD, SPF_PASS autolearn=ham version=3.3.2 spammy=b3, a1, b0, a2 X-HELO: fencepost.gnu.org Received: from fencepost.gnu.org (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Mon, 14 Mar 2016 23:18:24 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:55761) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1afbkk-000485-7H for gcc-patches@gnu.org; Mon, 14 Mar 2016 19:18:22 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1afbkg-0005rq-E7 for gcc-patches@gnu.org; Mon, 14 Mar 2016 19:18:21 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:41814) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1afbkg-0005rj-4i for gcc-patches@gnu.org; Mon, 14 Mar 2016 19:18:18 -0400 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1afbkc-0004b7-Ak from Tom_deVries@mentor.com ; Mon, 14 Mar 2016 16:18:14 -0700 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Mon, 14 Mar 2016 23:18:12 +0000 Subject: Re: [PATCH, 4/16] Implement -foffload-alias To: Jakub Jelinek References: <20151111110034.GF5675@tucnak.redhat.com> <5644B84D.6050504@mentor.com> <5645C33B.9080802@mentor.com> <20151113113938.GM5675@tucnak.redhat.com> <565058F0.8040509@mentor.com> <56584191.60704@mentor.com> <565846A8.6000509@mentor.com> <20151202095849.GE5675@tucnak.redhat.com> <56E6B99F.60507@mentor.com> CC: Richard Biener , "gcc-patches@gnu.org" From: Tom de Vries Message-ID: <56E746A2.3000207@mentor.com> Date: Tue, 15 Mar 2016 00:17:54 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.6.0 MIME-Version: 1.0 In-Reply-To: <56E6B99F.60507@mentor.com> X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 On 14/03/16 14:16, Tom de Vries wrote: > On 02/12/15 10:58, Jakub Jelinek wrote: >> On Fri, Nov 27, 2015 at 01:03:52PM +0100, Tom de Vries wrote: >>> Handle non-declared variables in kernels alias analysis >>> >>> 2015-11-27 Tom de Vries >>> >>> * gimplify.c (gimplify_scan_omp_clauses): Initialize >>> OMP_CLAUSE_ORIG_DECL. >>> * omp-low.c (install_var_field_1): Handle base_pointers_restrict for >>> pointers. >>> (map_ptr_clause_points_to_clause_p) >>> (nr_map_ptr_clauses_pointing_to_clause): New function. >>> (omp_target_base_pointers_restrict_p): Handle GOMP_MAP_POINTER. >>> * tree-pretty-print.c (dump_omp_clause): Print OMP_CLAUSE_ORIG_DECL. >>> * tree.c (omp_clause_num_ops): Set num_ops for OMP_CLAUSE_MAP to 3. >>> * tree.h (OMP_CLAUSE_ORIG_DECL): New macro. >>> >>> * c-c++-common/goacc/kernels-alias-10.c: New test. >>> * c-c++-common/goacc/kernels-alias-9.c: New test. >> >> I don't like this (mainly the addition of OMP_CLAUSE_ORIG_DECL), >> but it also sounds wrong to me. >> The primary question is how do you handle GOMP_MAP_POINTER >> (which is something we don't use for C/C++ OpenMP anymore, >> and Fortran OpenMP will stop using it in GCC 7 or 6.2?) on the OpenACC >> libgomp side, does it work like GOMP_MAP_ALLOC or GOMP_MAP_FORCE_ALLOC? > > When a GOMP_MAP_POINTER mapping is encountered, first we check if it has > been mapped before: > - if it hasn't been mapped before, we check if the area the pointer > points to has been mapped, and if not, error out. Else we map the > pointer to a device pointer, and write the device pointer value > to the device pointer variable. > - if the pointer has been mapped before, we reuse the mapping and write > the device pointer value to the device pointer variable. > >> Similarly GOMP_MAP_TO_PSET. >> If it works like GOMP_MAP_ALLOC (it does >> on the OpenMP side in target.c, so if something is already mapped, no >> further pointer assignment happens), then your change looks wrong. >> If it works like GOMP_MAP_FORCE_ALLOC, then you just should treat >> GOMP_MAP_POINTER on all OpenACC constructs as opcode that allows the >> restrict operation. > > I guess it works mostly like GOMP_MAP_ALLOC, but I don't understand the > relevance of the comparison for the patch. What is interesting for the > restrict optimization is whether what GOMP_MAP_POINTER points to has > been mapped with or without the force flag during the same mapping > sequence. > >> If it should behave differently depending on >> if the corresponding array section has been mapped with GOMP_MAP_FORCE_* >> or without it, > > The mapping itself shouldn't behave differently. > >> then supposedly you should use a different code for >> those two. > > I could add f.i. an unsigned int aux_flags to struct tree_omp_clause, > set a new POINTS_TO_FORCE_VAR flag when translating the acc clause into > mapping clauses, and use that flag later on when dealing with the > GOMP_MAP_POINTER clause. Is that an acceptable approach? > > [ Instead I could define a new gcc-internal-only > GOMP_MAP_POINTER_POINTS_TO_FORCE kind, but I'd rather avoid this, given > that it would be handled the same as GOMP_MAP_POINTER everywhere, except > for a single point in the source code. ] I found the example of OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION and OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION, which re-purpose existing but unused fields, and used something similar in attached patch (untested, c-only for the moment). Thanks, - Tom 2016-03-14 Tom de Vries * omp-low.c (install_var_field): Handle base_pointers_restrict for pointers. (omp_target_base_pointers_restrict_p): Handle GOMP_MAP_POINTER. * tree.h (OMP_CLAUSE_MAP_POINTER_TO_FORCED): define. * c-typeck.c (handle_omp_array_sections): Set OMP_CLAUSE_MAP_POINTER_TO_FORCED on GOMP_MAP_POINTER clause. * c-c++-common/goacc/kernels-alias-10.c: New test. * c-c++-common/goacc/kernels-alias-9.c: New test. Handle non-declared variables in kernels alias analysis --- gcc/c/c-typeck.c | 15 ++++++- gcc/omp-low.c | 48 ++++++++++++++++++++++ .../c-c++-common/goacc/kernels-alias-10.c | 29 +++++++++++++ gcc/testsuite/c-c++-common/goacc/kernels-alias-9.c | 29 +++++++++++++ gcc/tree.h | 3 ++ 5 files changed, 123 insertions(+), 1 deletion(-) diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 6aa0f03..a05831d 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12446,7 +12446,20 @@ handle_omp_array_sections (tree c, bool is_omp) } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (!is_omp) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); + { + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + OMP_CLAUSE_MAP_POINTER_TO_FORCED (c2) = 1; + break; + default: + break; + } + } else if (TREE_CODE (t) == COMPONENT_REF) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); else diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 82dec9d..f9d953d 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1429,6 +1429,9 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, } else if (by_ref) { + if (base_pointers_restrict + && POINTER_TYPE_P (type)) + type = build_qualified_type (type, TYPE_QUAL_RESTRICT); type = build_pointer_type (type); if (base_pointers_restrict) type = build_qualified_type (type, TYPE_QUAL_RESTRICT); @@ -3132,6 +3135,47 @@ omp_target_base_pointers_restrict_p (tree clauses) Because both mappings have the force prefix, we know that they will be allocated when calling the corresponding offloaded function, which means we can mark the base pointers for a and b in the offloaded function as + restrict. + + II. GOMP_MAP_POINTER example: + + void foo (unsigned int *a, unsigned int *b) + { + #pragma acc kernels copyout (a[0:2]) copyout (b[0:2]) + { + a[0] = 0; + b[0] = 1; + } + } + + After gimplification, we have: + + foo (unsigned int * a, unsigned int * b) + { + unsigned int * b.0; + unsigned int * a.1; + + b.0 = b; + a.1 = a; + #pragma omp target oacc_kernels \ + map(force_from:*a.1 (*a) [len: 8]) \ + map(alloc:a [pointer assign, bias: 0]) \ + map(force_from:*b.0 (*b) [len: 8]) \ + map(alloc:b [pointer assign, bias: 0]) + { + unsigned int * a.2; + unsigned int * b.3; + + a.2 = a; + *a.2 = 0; + b.3 = b; + *b.3 = 1; + } + } + + By testing for OMP_CLAUSE_MAP_POINTER_TO_FORCED, we can known for both + pointer assign mappings that they point to a force-prefixed mapping, so + we can mark the base pointers for a and b in the offloaded function as restrict. */ tree c; @@ -3147,6 +3191,10 @@ omp_target_base_pointers_restrict_p (tree clauses) case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TOFROM: break; + case GOMP_MAP_POINTER: + if (!OMP_CLAUSE_MAP_POINTER_TO_FORCED (c)) + return false; + break; default: return false; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-10.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-10.c new file mode 100644 index 0000000..ce5bbe8 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-10.c @@ -0,0 +1,29 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + unsigned int b[N]; + unsigned int c[N]; + unsigned int d[N]; + +#pragma acc kernels copyin (a[0:N]) create (b[0:N]) copyout (c[0:N]) copy (d[0:N]) + { + a[0] = 0; + b[0] = 0; + c[0] = 0; + d[0] = 0; + } +} + +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-9.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-9.c new file mode 100644 index 0000000..7229fd4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-9.c @@ -0,0 +1,29 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +#define N 2 + +void +foo (unsigned int *a, unsigned int *b, unsigned int *c, unsigned int *d) +{ + +#pragma acc kernels copyin (a[0:N]) create (b[0:N]) copyout (c[0:N]) copy (d[0:N]) + { + a[0] = 0; + b[0] = 0; + c[0] = 0; + d[0] = 0; + } +} + +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 6" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 7" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 8" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 9" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 12 "ealias" } } */ + diff --git a/gcc/tree.h b/gcc/tree.h index 544a6a1..bc48ea8 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1533,6 +1533,9 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \ TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) +#define OMP_CLAUSE_MAP_POINTER_TO_FORCED(NODE) \ + TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) + #define OMP_CLAUSE_PROC_BIND_KIND(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)