[Openmp-commits] [PATCH] D110193: [RFC] Initial documentation for declare target indirect support.

[Deepak Eachempati via Phabricator via Openmp-commits]

dreachem added inline comments.


================
Comment at: openmp/libomptarget/docs/declare_target_indirect.md:39
+
+Since all pointers are supposed to be translated/mapped, when program does not use **required unified_shared_memory**, it is possible to avoid generating the runtime dispatch code for indirect function calls. The mapping between host and device address of an indirect function will be established by `libomptarget` during processing of the offload entries table.
+
----------------
vzakhari wrote:
> RaviNarayanaswamy wrote:
> > vzakhari wrote:
> > > jdoerfert wrote:
> > > > How would we translate/map all pointers? I fail to see how we could even identify all (function) pointers (see below). Or did we add wording to the standard to forbid this?
> > > > 
> > > > ```
> > > > char data[sizeof(fp)];
> > > > memcpy(data, &fp, sizeof(fp));
> > > > 
> > > > #pragma omp target map(to:data[:sizeof(fp)])
> > > > {
> > > >   void (*fpc)();
> > > >   memcpy(&fpc, data, sizeof(fp));
> > > >   fpc(); 
> > > > }
> > > > ```
> > > > 
> > > > Modulo my syntax errors ;)
> > > I agree that in this case there will be no translation.  I will remove the words about non-unified_shared_memory "optimization".  Does it sound right?
> > > I agree that in this case there will be no translation.  I will remove the words about non-unified_shared_memory "optimization".  Does it sound right?
> > 
> > We should get the spec to disallow copy/casting of host pointers and using it  to allow the optimization.
> I can move this to a TODO section at the end of the doc.  Let's agree on the general case, first.
> > I agree that in this case there will be no translation.  I will remove the words about non-unified_shared_memory "optimization".  Does it sound right?
> 
> We should get the spec to disallow copy/casting of host pointers and using it  to allow the optimization.

This restriction was added for 5.2:  "A program must not rely on the value of a function address in a **target** region except for assignments, comparisons to zero and indirect calls."

So, you can copy pointers around, as in the example posted. You want to forbid this as well?


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D110193/new/

https://reviews.llvm.org/D110193