[Openmp-commits] [PATCH] D117997: [libomptarget] Lookup by base address when begin address is not mapped.

Ye Luo via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Sun Jan 23 14:02:20 PST 2022

ye-luo added inline comments.

Comment at: openmp/libomptarget/test/mapping/array_section_implicit_capture.c:31
+  // are captured implicitly so the library must look them up using their base
+  // address.
+  #pragma omp target
grokos wrote:
> ABataev wrote:
> > grokos wrote:
> > > ye-luo wrote:
> > > > When two distinct sections of A is mapped, implicit map looks up the base address, then which section to pick.
> > > > 
> > > > I think in this example, the implicit map should fail to look up A.
> > > > Users needs to explicitly add map(to: A[FROM:LENGTH], B[FROM:LENGTH]) to be valid.
> > > We only map one section, mapping two distinct sections of the same object is illegal. Implicitly mapping more, i.e. asking for the base when only a section in the middle of the object has been mapped is allowed. This example is valid OpenMP usage.
> > IIRC, Ye is correct here. Mapping 2 sections in different constructs is legal, just the inner one must be a subset of the outer.
> Sure, re-mapping what is already mapped is perfectly fine. I'm talking about implicitly "mapping" more, i.e. not really mapping a superset but rather using a zero-length pointer (or the whole named array) to refer to the object just for the sake of convenience (without having to re-write the exact range). Libomptarget has always supported this scenario, e.g. we have mapped A[0 : N/2] and then pass A[0:0] (or A[0:N] if A is a named array) to the target region. Now, if we split the array across multiple devices and each device only maps a chunk, without this patch all devices (except for the first one) will fail to look up the mapping.
I think the spec has taken into account what I have mentioned and there is `5.8.6 Pointer Initialization for Device Data Environments` in the spec. Clearly, base pointer is the way to go.

  rG LLVM Github Monorepo



More information about the Openmp-commits mailing list