From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 3719 invoked by alias); 26 Oct 2015 18:34:31 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 3708 invoked by uid 89); 26 Oct 2015 18:34:30 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Mon, 26 Oct 2015 18:34:29 +0000 Received: from int-mx13.intmail.prod.int.phx2.redhat.com (int-mx13.intmail.prod.int.phx2.redhat.com [10.5.11.26]) by mx1.redhat.com (Postfix) with ESMTPS id 5BE26A2C3E; Mon, 26 Oct 2015 18:34:28 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-121.ams2.redhat.com [10.36.116.121]) by int-mx13.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t9QIYQB8007766 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Mon, 26 Oct 2015 14:34:27 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id t9QIYO5P006543; Mon, 26 Oct 2015 19:34:25 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id t9QIYMf0006542; Mon, 26 Oct 2015 19:34:22 +0100 Date: Mon, 26 Oct 2015 18:36:00 -0000 From: Jakub Jelinek To: James Norris Cc: GCC Patches , "Joseph S. Myers" , Nathan Sidwell Subject: Re: [Bulk] [OpenACC 0/7] host_data construct Message-ID: <20151026183422.GW478@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <56293476.5020801@codesourcery.com> <562A578E.4080907@codesourcery.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <562A578E.4080907@codesourcery.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-10/txt/msg02799.txt.bz2 On Fri, Oct 23, 2015 at 10:51:42AM -0500, James Norris wrote: > @@ -12942,6 +12961,7 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) > case OMP_CLAUSE_GANG: > case OMP_CLAUSE_WORKER: > case OMP_CLAUSE_VECTOR: > + case OMP_CLAUSE_USE_DEVICE: > pc = &OMP_CLAUSE_CHAIN (c); > continue; > Are there any restrictions on whether you can specify the same var multiple times in use_device clause? #pragma acc host_data use_device (x) use_device (x) use_device (y, y, y) ? If not, have you verified that the gimplifier doesn't ICE on it? Generally it doesn't like the same var being mentioned multiple times. If yes, you can use e.g. the generic_head bitmap for that and in any case, cover that with sufficient testsuite coverage. > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index ab9e540..0c32219 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -93,6 +93,8 @@ enum gimplify_omp_var_data > > GOVD_MAP_0LEN_ARRAY = 32768, > > + GOVD_USE_DEVICE = 65536, > + > GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE > | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR > | GOVD_LOCAL) > @@ -116,7 +118,9 @@ enum omp_region_type > ORT_COMBINED_TARGET = 33, > /* Dummy OpenMP region, used to disable expansion of > DECL_VALUE_EXPRs in taskloop pre body. */ > - ORT_NONE = 64 > + ORT_NONE = 64, > + /* An OpenACC host-data region. */ > + ORT_HOST_DATA = 128 I'd prefer ORT_NONE to be the last one, can you just renumber it and put ORT_HOST_DATA before it? > +static tree > +gimplify_oacc_host_data_1 (tree *tp, int *walk_subtrees, > + void *data ATTRIBUTE_UNUSED) > +{ Your use_device sounds very similar to use_device_ptr clause in OpenMP, which is allowed on #pragma omp target data construct and is implemented quite a bit differently from this; it is unclear if the OpenACC standard requires this kind of implementation, or you just chose to implement it this way. In particular, the GOMP_target_data call puts the variables mentioned in the use_device_ptr clauses into the mapping structures (similarly how map clause appears) and the corresponding vars are privatized within the target data region (which is a host region, basically a fancy { } braces), where the private variables contain the offloading device's pointers. > + splay_tree_node n = NULL; > + location_t loc = EXPR_LOCATION (*tp); > + > + switch (TREE_CODE (*tp)) > + { > + case ADDR_EXPR: > + { > + tree decl = TREE_OPERAND (*tp, 0); > + > + switch (TREE_CODE (decl)) > + { > + case ARRAY_REF: > + case ARRAY_RANGE_REF: > + case COMPONENT_REF: > + case VIEW_CONVERT_EXPR: > + case REALPART_EXPR: > + case IMAGPART_EXPR: > + if (TREE_CODE (TREE_OPERAND (decl, 0)) == VAR_DECL) > + n = splay_tree_lookup (gimplify_omp_ctxp->variables, > + (splay_tree_key) TREE_OPERAND (decl, 0)); > + break; I must say this looks really strange, you throw away all the offsets embedded in the component codes (fixed or variable). Where comes the above list? What about other components (say bit field refs, etc.)? > + case VAR_DECL: What is so special about VAR_DECLs? Shouldn't PARM_DECLs / RESULT_DECLs be treated the same way? > --- a/libgomp/libgomp.map > +++ b/libgomp/libgomp.map > @@ -378,6 +378,7 @@ GOACC_2.0 { > GOACC_wait; > GOACC_get_thread_num; > GOACC_get_num_threads; > + GOACC_deviceptr; > }; > > GOACC_2.0.1 { You shouldn't be adding new symbols into a symbol version that appeared in a compiler that shipped already (GCC 5 already had GOACC_2.0 symbols). So it should go into GOACC_2.0.1. > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c > index af067d6..497ab92 100644 > --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > @@ -204,6 +204,38 @@ acc_deviceptr (void *h) > return d; > } > > +/* This function is used as a helper in generated code to implement pointer > + lookup in host_data regions. Unlike acc_deviceptr, it returns its argument > + unchanged on a shared-memory system (e.g. the host). */ > + > +void * > +GOACC_deviceptr (void *h) > +{ > + splay_tree_key n; > + void *d; > + void *offset; > + > + goacc_lazy_initialize (); > + > + struct goacc_thread *thr = goacc_thread (); > + > + if ((thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) == 0) > + { > + n = lookup_host (thr->dev, h, 1); What is supposed to be the behavior when the h pointer points at object boundary, rather than into the middle of existing mapped object? Say you have: char a[16], b[0], c[16]; // b is GCC extension Now, char *p = &a[5]; is unambiguous, either a is mapped, or not. But, if p = &a[16];, then it could be either the one-past-last byte in a, or it could be the start of b (== one-past-last byte in b) or it could be the pointer to start of c. In OpenMP 4.5, I had endless discussions about this and the end result is that one-past-last byte addresses are unspecified behavior Jakub