2015-08-14 11:47 GMT+02:00 Thomas Schwinge <tho...@codesourcery.com>:
> On Wed, 5 Aug 2015 18:09:04 +0300, Ilya Verbin <iver...@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
