From patchwork Fri Nov 27 11:42:09 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 549405 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 2A4291402A5 for ; Fri, 27 Nov 2015 22:44:00 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=v3Edemi4; 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=ArxN+NBWE2blXvixd aqM20/HwZUbqVqkaMZ+rZ9+6fkvm9/eGaJATcgqTvFwSQzkDC7wTwDd31hQwTHxQ QTdVZqvi98xiSrHBAtbxm9l6xWOwK4h440tH7wNmHF6loTnLwq9gu5NT0n+YLqZ0 snqM8eT2Xd+cRReh53QMvMSW58= 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=sDPmYwG7/s7ucea/1/yNkuq mL04=; b=v3Edemi45b8Sz2urhlRiagts4khiQKUsii/5VYEMgwo/VDVXm4XeKK4 OLmfSDEl0tE8f8dgFEnlpXqfICZRmfrjaBi+5oI7UnECJ3++iOm4bXJ19e91DJUm i4kfNXFLkSn/bH/C/gu5TzDmqXN2LBprCpl9ZraGUP82G1Hg1DhA= Received: (qmail 114884 invoked by alias); 27 Nov 2015 11:43:38 -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 114744 invoked by uid 89); 27 Nov 2015 11:43:36 -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, SPF_PASS, T_RP_MATCHES_RCVD autolearn=ham version=3.3.2 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; Fri, 27 Nov 2015 11:43:16 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:44792) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1a2HQm-0000Pc-UG for gcc-patches@gnu.org; Fri, 27 Nov 2015 06:43:13 -0500 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1a2HQj-0007Nf-BD for gcc-patches@gnu.org; Fri, 27 Nov 2015 06:43:12 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:39287) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1a2HQj-0007Ml-0v for gcc-patches@gnu.org; Fri, 27 Nov 2015 06:43:09 -0500 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 1a2HQf-0003Aj-BV from Tom_deVries@mentor.com ; Fri, 27 Nov 2015 03:43:06 -0800 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; Fri, 27 Nov 2015 11:43:03 +0000 Subject: Re: [PATCH, 4/16] Implement -foffload-alias To: Richard Biener References: <5640BD31.2060602@mentor.com> <5640C560.1000007@mentor.com> <20151111110034.GF5675@tucnak.redhat.com> <5644B84D.6050504@mentor.com> <5645C33B.9080802@mentor.com> <20151113113938.GM5675@tucnak.redhat.com> <565058F0.8040509@mentor.com> CC: Jakub Jelinek , "gcc-patches@gnu.org" From: Tom de Vries Message-ID: <56584191.60704@mentor.com> Date: Fri, 27 Nov 2015 12:42:09 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 On 23/11/15 12:41, Richard Biener wrote: > On Sat, 21 Nov 2015, Tom de Vries wrote: > >> >On 13/11/15 12:39, Jakub Jelinek wrote: >>> > >On Fri, Nov 13, 2015 at 12:29:51PM +0100, Richard Biener wrote: >>>>> > > > >thanks for the explanation. Filed as PR68331 - '[meta-bug] fipa-pta >>>>> > > > >issues'. >>>>> > > > > >>>>> > > > >Any feedback on the '#pragma GCC offload-alias=' bit >>>>> > > > >above? >>>>> > > > >Is that sort of what you had in mind? >>>> > > > >>>> > > >Yes. Whether that makes sense is another question of course. You can >>>> > > >annotate memory references with MR_DEPENDENCE_BASE/CLIQUE yourself >>>> > > >as well if you know dependences without the users intervention. >>> > > >>> > >I really don't like even the GCC offload-alias, I just don't see anything >>> > >special on the offload code. Not to mention that the same issue is already >>> > >with other outlined functions, like OpenMP tasks or parallel regions, those >>> > >aren't offloaded, yet they can suffer from worse alias/points-to analysis >>> > >too. >> > >> >AFAIU there is one aspect that is different for offloaded code: the setup of >> >the data on the device. >> > >> >Consider this example: >> >... >> >unsigned int a[N]; >> >unsigned int b[N]; >> >unsigned int c[N]; >> > >> >int >> >main (void) >> >{ >> > ... >> > >> >#pragma acc kernels copyin (a) copyin (b) copyout (c) >> > { >> > for (COUNTERTYPE ii = 0; ii < N; ii++) >> > c[ii] = a[ii] + b[ii]; >> > } >> > >> > ... >> >... >> > >> >At gimple level, we have: >> >... >> >#pragma omp target oacc_kernels \ >> > map(force_from:c [len: 2097152]) \ >> > map(force_to:b [len: 2097152]) \ >> > map(force_to:a [len: 2097152]) >> >... >> > >> >[ The meaning of the force_from/force_to mappings is given in >> >include/gomp-constants.h: >> >... >> > /* Allocate. */ >> > GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC), >> > /* ..., and copy to device. */ >> > GOMP_MAP_FORCE_TO = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TO), >> > /* ..., and copy from device. */ >> > GOMP_MAP_FORCE_FROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM), >> > /* ..., and copy to and from device. */ >> > GOMP_MAP_FORCE_TOFROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM), >> >... ] >> > >> >So before calling the offloaded function, a separate alloc is done for a, b >> >and c, and the base pointers of the newly allocated objects are passed to the >> >offloaded function. >> > >> >This means we can mark those base pointers as restrict in the offloaded >> >function. >> > >> >Attached proof-of-concept patch implements that. >> > >>> > >We simply have some compiler internal interface between the caller and >>> > >callee of the outlined regions, each interface in between those has >>> > >its own structure type used to communicate the info; >>> > >we can attach attributes on the fields, or some flags to indicate some >>> > >properties interesting from aliasing POV. >>> > >We don't really need to perform >>> > >full IPA-PTA, perhaps it would be enough to a) record somewhere in cgraph >>> > >the relationship in between such callers and callees (for offloading regions >>> > >we already have "omp target entrypoint" attribute on the callee and a >>> > >singler caller), tell LTO if possible not to split those into different >>> > >partitions if easily possible, and then just for these pairs perform >>> > >aliasing/points-to analysis in the caller and the result record using >>> > >cliques/special attributes/whatever to the callee side, so that the callee >>> > >(outlined OpenMP/OpenACC/Cilk+ region) can then improve its alias analysis. >> > >> >As a start, is the approach of this patch OK? > Works for me but leaving to Jakub to review for correctness. Attached patch is a complete version: - added ChangeLog - added missing function header comments - moved analysis to separate function omp_target_base_pointers_restrict_p - added example in comment before analysis - fixed error in omp_target_base_pointers_restrict_p where I was using GOMP_MAP_ALLOC but should have been using GOMP_MAP_FORCE_ALLOC - added testcases Bootstrapped and reg-tested on x86_64. OK for stage3 trunk? Thanks, - Tom Mark pointers to allocated target vars as restricted, if possible 2015-11-26 Tom de Vries * omp-low.c (install_var_field_1): New function, factored out of ... (install_var_field): ... here. (scan_sharing_clauses_1): New function, factored out of ... (scan_sharing_clauses): ... here. (omp_target_base_pointers_restrict_p): New function. (scan_omp_target): Call scan_sharing_clauses_1 instead of scan_sharing_clauses, with base_pointers_restrict arg. * c-c++-common/goacc/kernels-alias-2.c: New test. * c-c++-common/goacc/kernels-alias-3.c: New test. * c-c++-common/goacc/kernels-alias-4.c: New test. * c-c++-common/goacc/kernels-alias-5.c: New test. * c-c++-common/goacc/kernels-alias-6.c: New test. * c-c++-common/goacc/kernels-alias-7.c: New test. * c-c++-common/goacc/kernels-alias-8.c: New test. * c-c++-common/goacc/kernels-alias.c: New test. --- gcc/omp-low.c | 109 +++++++++++++++++++-- gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c | 27 +++++ gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c | 20 ++++ gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c | 22 +++++ gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c | 19 ++++ gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c | 23 +++++ gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c | 25 +++++ gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c | 22 +++++ gcc/testsuite/c-c++-common/goacc/kernels-alias.c | 29 ++++++ 9 files changed, 289 insertions(+), 7 deletions(-) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 0d4c6e5..6843c49 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1366,10 +1366,12 @@ build_sender_ref (tree var, omp_context *ctx) return build_sender_ref ((splay_tree_key) var, ctx); } -/* Add a new field for VAR inside the structure CTX->SENDER_DECL. */ +/* Add a new field for VAR inside the structure CTX->SENDER_DECL. If + BASE_POINTERS_RESTRICT, declare the field with restrict. */ static void -install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) +install_var_field_1 (tree var, bool by_ref, int mask, omp_context *ctx, + bool base_pointers_restrict) { tree field, type, sfield = NULL_TREE; splay_tree_key key = (splay_tree_key) var; @@ -1393,7 +1395,11 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) type = build_pointer_type (build_pointer_type (type)); } else if (by_ref) - type = build_pointer_type (type); + { + type = build_pointer_type (type); + if (base_pointers_restrict) + type = build_qualified_type (type, TYPE_QUAL_RESTRICT); + } else if ((mask & 3) == 1 && is_reference (var)) type = TREE_TYPE (type); @@ -1457,6 +1463,14 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) splay_tree_insert (ctx->sfield_map, key, (splay_tree_value) sfield); } +/* As install_var_field_1, but with base_pointers_restrict == false. */ + +static void +install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) +{ + install_var_field_1 (var, by_ref, mask, ctx, false); +} + static tree install_var_local (tree var, omp_context *ctx) { @@ -1810,10 +1824,12 @@ fixup_child_record_type (omp_context *ctx) } /* Instantiate decls as necessary in CTX to satisfy the data sharing - specified by CLAUSES. */ + specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with + restrict. */ static void -scan_sharing_clauses (tree clauses, omp_context *ctx) +scan_sharing_clauses_1 (tree clauses, omp_context *ctx, + bool base_pointers_restrict) { tree c, decl; bool scan_array_reductions = false; @@ -2070,7 +2086,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) install_var_field (decl, true, 7, ctx); else - install_var_field (decl, true, 3, ctx); + install_var_field_1 (decl, true, 3, ctx, + base_pointers_restrict); if (is_gimple_omp_offloaded (ctx->stmt)) install_var_local (decl, ctx); } @@ -2336,6 +2353,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) scan_omp (&OMP_CLAUSE_LINEAR_GIMPLE_SEQ (c), ctx); } +/* As scan_sharing_clauses_1, but with base_pointers_restrict == false. */ + +static void +scan_sharing_clauses (tree clauses, omp_context *ctx) +{ + scan_sharing_clauses_1 (clauses, ctx, false); +} + /* Create a new name for omp child function. Returns an identifier. If IS_CILK_FOR is true then the suffix for the child function is "_cilk_for_fn." */ @@ -3032,6 +3057,68 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx) layout_type (ctx->record_type); } +/* Return true if the CLAUSES of an omp target guarantee that the base pointers + used in the corresponding offloaded function are restrict. */ + +static bool +omp_target_base_pointers_restrict_p (tree clauses) +{ + /* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only + used by OpenACC. */ + if (flag_openacc == 0) + return false; + + /* I. Basic example: + + void foo (void) + { + unsigned int a[2], b[2]; + + #pragma acc kernels \ + copyout (a) \ + copyout (b) + { + a[0] = 0; + b[0] = 1; + } + } + + After gimplification, we have: + + #pragma omp target oacc_kernels \ + map(force_from:a [len: 8]) \ + map(force_from:b [len: 8]) + { + a[0] = 0; + b[0] = 1; + } + + 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. */ + + tree c; + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + return false; + + 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: + break; + default: + return false; + } + } + + return true; +} + /* Scan a GIMPLE_OMP_TARGET. */ static void @@ -3053,13 +3140,21 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) DECL_NAMELESS (name) = 1; TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; + + bool base_pointers_restrict = false; if (offloaded) { create_omp_child_function (ctx, false); gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn); + + base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses); + if (base_pointers_restrict + && dump_file && (dump_flags & TDF_DETAILS)) + fprintf (dump_file, + "Base pointers in offloaded function are restrict\n"); } - scan_sharing_clauses (clauses, ctx); + scan_sharing_clauses_1 (clauses, ctx, base_pointers_restrict); scan_omp (gimple_omp_body_ptr (stmt), ctx); if (TYPE_FIELDS (ctx->record_type) == NULL) diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c new file mode 100644 index 0000000..d437c47 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +void +foo (void) +{ + unsigned int a; + unsigned int b; + unsigned int c; + unsigned int d; + +#pragma acc kernels copyin (a) create (b) copyout (c) copy (d) + { + a = 0; + b = 0; + c = 0; + d = 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-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c new file mode 100644 index 0000000..0eda7e1 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c @@ -0,0 +1,20 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +void +foo (void) +{ + unsigned int a; + unsigned int *p = &a; + +#pragma acc kernels pcopyin (a, p[0:1]) + { + a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c new file mode 100644 index 0000000..037901f --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c @@ -0,0 +1,22 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + unsigned int *p = &a[0]; + +#pragma acc kernels pcopyin (a, p[0:2]) + { + a[0] = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c new file mode 100644 index 0000000..69cd3fb --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c @@ -0,0 +1,19 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +void +foo (int *a) +{ + int *p = a; + +#pragma acc kernels pcopyin (a[0:1], p[0:1]) + { + *a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c new file mode 100644 index 0000000..6ebce15 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c @@ -0,0 +1,23 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +typedef __SIZE_TYPE__ size_t; +extern void *acc_copyin (void *, size_t); + +void +foo (void) +{ + int a = 0; + int *p = (int *)acc_copyin (&a, sizeof (a)); + +#pragma acc kernels deviceptr (p) pcopy(a) + { + a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c new file mode 100644 index 0000000..40eb235 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c @@ -0,0 +1,25 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +typedef __SIZE_TYPE__ size_t; +extern void *acc_copyin (void *, size_t); + +#define N 2 + +void +foo (void) +{ + int a[N]; + int *p = (int *)acc_copyin (&a[0], sizeof (a)); + +#pragma acc kernels deviceptr (p) pcopy(a) + { + a[0] = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c new file mode 100644 index 0000000..0b93e35 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c @@ -0,0 +1,22 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +typedef __SIZE_TYPE__ size_t; +extern void *acc_copyin (void *, size_t); + +void +foo (int *a, size_t n) +{ + int *p = (int *)acc_copyin (&a, n); + +#pragma acc kernels deviceptr (p) pcopy(a[0:n]) + { + a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c new file mode 100644 index 0000000..25821ab2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias.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) create (b) copyout (c) copy (d) + { + 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" } } */ +