This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: Forwarding -foffload=[...] from the driver (compile-time) to libgomp (run-time) (was: [PATCH 2/n] OpenMP 4.0 offloading infrastructure: LTO streaming)
- From: Ilya Verbin <iverbin at gmail dot com>
- To: Thomas Schwinge <thomas at codesourcery dot com>
- Cc: Jakub Jelinek <jakub at redhat dot com>, Richard Biener <richard dot guenther at gmail dot com>, Joseph Myers <joseph at codesourcery dot com>, Richard Biener <rguenther at suse dot de>, Jan Hubicka <hubicka at ucw dot cz>, GCC Patches <gcc-patches at gcc dot gnu dot org>, Kirill Yukhin <kirill dot yukhin at gmail dot com>, mjambor at suse dot cz
- Date: Fri, 14 Aug 2015 15:19:26 +0200
- Subject: Re: Forwarding -foffload=[...] from the driver (compile-time) to libgomp (run-time) (was: [PATCH 2/n] OpenMP 4.0 offloading infrastructure: LTO streaming)
- Authentication-results: sourceware.org; auth=none
- References: <alpine dot LSU dot 2 dot 11 dot 1410151612470 dot 20733 at zhemvz dot fhfr dot qr> <20141020111935 dot GA9362 at msticlxl57 dot ims dot intel dot com> <20141024141601 dot GA62562 at msticlxl57 dot ims dot intel dot com> <20141024142028 dot GD10376 at tucnak dot redhat dot com> <20141028193047 dot GA17865 at msticlxl57 dot ims dot intel dot com> <20141103092447 dot GO5026 at tucnak dot redhat dot com> <20141105124655 dot GA42356 at msticlxl57 dot ims dot intel dot com> <87egjopgh0 dot fsf at kepler dot schwinge dot homeip dot net> <20150731142007 dot GA64740 at msticlxl57 dot ims dot intel dot com> <CAFiYyc3NJjund2BbEQ0HxtGiaM-Z9bmX4UPaQtYguYjHQs0-Gg at mail dot gmail dot com> <20150805150904 dot GA3211 at msticlxl57 dot ims dot intel dot com> <87bneatd5q dot fsf at schwinge dot name>
2015-08-14 11:47 GMT+02:00 Thomas Schwinge <thomas@codesourcery.com>:
> On Wed, 5 Aug 2015 18:09:04 +0300, Ilya Verbin <iverbin@gmail.com> 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.
* 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