From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 116526 invoked by alias); 17 Aug 2015 13:53:30 -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 116513 invoked by uid 89); 17 Aug 2015 13:53:29 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.9 required=5.0 tests=AWL,BAYES_00,SPF_PASS autolearn=ham version=3.3.2 X-HELO: mx2.suse.de Received: from mx2.suse.de (HELO mx2.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (CAMELLIA256-SHA encrypted) ESMTPS; Mon, 17 Aug 2015 13:53:28 +0000 Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 53842AD9E; Mon, 17 Aug 2015 13:53:25 +0000 (UTC) Date: Mon, 17 Aug 2015 13:57:00 -0000 From: Martin Jambor To: Ilya Verbin Cc: Thomas Schwinge , Jakub Jelinek , Richard Biener , Joseph Myers , Richard Biener , Jan Hubicka , GCC Patches , Kirill Yukhin Subject: Re: Forwarding -foffload=[...] from the driver (compile-time) to libgomp (run-time) (was: [PATCH 2/n] OpenMP 4.0 offloading infrastructure: LTO streaming) Message-ID: <20150817135324.GZ2569@virgil.suse.cz> Mail-Followup-To: Ilya Verbin , Thomas Schwinge , Jakub Jelinek , Richard Biener , Joseph Myers , Richard Biener , Jan Hubicka , GCC Patches , Kirill Yukhin References: <20141024142028.GD10376@tucnak.redhat.com> <20141028193047.GA17865@msticlxl57.ims.intel.com> <20141103092447.GO5026@tucnak.redhat.com> <20141105124655.GA42356@msticlxl57.ims.intel.com> <87egjopgh0.fsf@kepler.schwinge.homeip.net> <20150731142007.GA64740@msticlxl57.ims.intel.com> <20150805150904.GA3211@msticlxl57.ims.intel.com> <87bneatd5q.fsf@schwinge.name> MIME-Version: 1.0 Content-Type: text/plain; charset=utf-8 Content-Disposition: inline In-Reply-To: User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-08/txt/msg00904.txt.bz2 Hi, On Fri, Aug 14, 2015 at 03:19:26PM +0200, Ilya Verbin wrote: > 2015-08-14 11:47 GMT+02:00 Thomas Schwinge : > > On Wed, 5 Aug 2015 18:09:04 +0300, Ilya Verbin wrote: > >> > > @@ -1095,6 +1092,8 @@ GOMP_target (int device, void (*fn) (void *), const void *unused, > >> > > return gomp_target_fallback (fn, hostaddrs); > >> > > > >> > > void *fn_addr = gomp_get_target_fn_addr (devicep, fn); > >> > > + if (fn_addr == NULL) > >> > > + return gomp_target_fallback (fn, hostaddrs); > > > > Is that reliable? Consider the following scenario, with f1 and f2 > > implemented in separate TUs: > > > > #pragma omp target data [map clauses] > > { > > f1([...]); > > f2([...]); > > } > > > > Consider that in f1 we have a OpenMP target region with offloading data > > available, and in f2 we have a OpenMP target region without offloading > > data available. In this case, the GOMP_target in f1 will execute on the > > offloading target, but the GOMP_target in f2 will resort to host fallback > > -- and we then likely have data inconsistencies, as the data specified by > > the map clauses is not synchronized between host and device. > > > > Admittedly, this is user error (inconsistent set of offloading functions > > available -- need either all, or none), but in such a scenario probably > > we should be doing a better job (at detecting this). (Note, I'm not sure > > whether my current patch actually does any better.) ;-) > > You're right. That's why I didn't send this patch for review yet. > My current plan is as follows: > * Use this approach for architectures with shared memory, since it > allows mixing host and target functions. Great, please keep me posted on these changes. Thanks! Martin > * For non-shared memory, at the first splay tree lookup: > ** If target fn is not found, run the whole program in host-fallback mode. > ** If it's found, then all target fns must exist. I.e. if some > tgt_addr (not first) is NULL, then libgomp will issue an error as it > does now. > > -- Ilya