Page MenuHomePhabricator

[RFC] Initial documentation for declare target indirect support.
ClosedPublic

Authored by vzakhari on Sep 21 2021, 12:25 PM.

Details

Summary

Please review proposed implementation of the support for declare target indirect.

Diff Detail

Event Timeline

vzakhari requested review of this revision.Sep 21 2021, 12:25 PM
vzakhari created this revision.
Herald added a project: Restricted Project. · View Herald Transcript
vzakhari added inline comments.Sep 21 2021, 12:36 PM
openmp/libomptarget/docs/declare_target_indirect.md
35

If needed, we can isolate the lookup implementation into a device RTL API, so that FE just generates a lookup call to map the original function address to the device function address.

jdoerfert added inline comments.Sep 21 2021, 6:35 PM
openmp/libomptarget/docs/declare_target_indirect.md
35

API, please: fp() -> (lookupDevPtr(fp))()

39

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 ;)

vzakhari added inline comments.Sep 21 2021, 7:03 PM
openmp/libomptarget/docs/declare_target_indirect.md
35

Okay, I will add a definition for the lookup API.

39

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?

openmp/libomptarget/docs/declare_target_indirect.md
39

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.

vzakhari added inline comments.Sep 22 2021, 10:13 AM
openmp/libomptarget/docs/declare_target_indirect.md
35

If I just specify the device RTL API that FEs will be able to use, I think I should remove any mentioning of __openmp_offload_function_ptr_map symbol. This becomes an implementation detail of the API for a particular device, so different device compiler and device plugins will be able to implement this different ways.

I can describe __openmp_offload_function_ptr_map as one of the alternative implementations. Does it sound right?

39

I can move this to a TODO section at the end of the doc. Let's agree on the general case, first.

jdoerfert added inline comments.Sep 22 2021, 10:43 AM
openmp/libomptarget/docs/declare_target_indirect.md
35

sounds good to me.

vzakhari updated this revision to Diff 374368.Sep 22 2021, 2:25 PM
vzakhari marked an inline comment as done.
dreachem added inline comments.Oct 3 2021, 4:33 PM
openmp/libomptarget/docs/declare_target_indirect.md
39

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?

vzakhari added inline comments.Oct 5 2021, 10:18 AM
openmp/libomptarget/docs/declare_target_indirect.md
39

I think Ravi only talked about disallowing verbatim copies of function addresses at the boundary of target regions, which will allow the optimization. @RaviNarayanaswamy, can you please reply to the comment?

openmp/libomptarget/docs/declare_target_indirect.md
39

Correct. Don't was passing function pointers as non function pointers in map clause. Assigning function pointer to another function pointer should be ok.

If there are no more comments, I am going to merge this by the end of the week.

JonChesterfield accepted this revision.Nov 8 2021, 3:00 AM

Let's have the docs change. They can always be updated during implementation.

This revision is now accepted and ready to land.Nov 8 2021, 3:00 AM