Hi,
On Fri, Aug 14, 2015 at 03:19:26PM +0200, Ilya Verbin wrote:
> 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.
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
