From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 46487 invoked by alias); 27 Oct 2015 15:45:36 -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 46478 invoked by uid 89); 27 Oct 2015 15:45:35 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.1 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 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; Tue, 27 Oct 2015 15:45:34 +0000 Received: from svr-orw-fem-06.mgc.mentorg.com ([147.34.97.120]) by relay1.mentorg.com with esmtp id 1Zr6RG-0002au-Kg from Cesar_Philippidis@mentor.com ; Tue, 27 Oct 2015 08:45:30 -0700 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-FEM-06.mgc.mentorg.com (147.34.97.120) with Microsoft SMTP Server id 14.3.224.2; Tue, 27 Oct 2015 08:45:30 -0700 Subject: Re: Re: [Bulk] [OpenACC 0/7] host_data construct To: Jakub Jelinek , James Norris References: <20151026183422.GW478@tucnak.redhat.com> CC: GCC Patches , "Joseph S. Myers" , Nathan Sidwell From: Cesar Philippidis Message-ID: <562F9C1A.7060905@codesourcery.com> Date: Tue, 27 Oct 2015 15:57:00 -0000 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: <20151026183422.GW478@tucnak.redhat.com> Content-Type: text/plain; charset="windows-1252" Content-Transfer-Encoding: 7bit X-SW-Source: 2015-10/txt/msg02925.txt.bz2 On 10/26/2015 11:34 AM, Jakub Jelinek wrote: > 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. Generally variables cannot appear in multiple clauses. I'll add more testing for this. >> 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? OK. >> +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. Is this a new OpenMP 4.5 feature? I'll take a closer look and see if they are similar enough. I also noticed that OpenMP 4.5 has something similar to OpenACC's enter/exit data construct now. >> + 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.)? I'm not sure. This is one of those things where multiple developers worked on it, and the history got lost. I'll investigate it. >> + 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. OK. >> 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? Probably undefined with the way that OpenACC is defined. > 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 OK. Thanks for you feedback. Cesar