[Openmp-commits] [PATCH] D117997: [libomptarget] Correctly return the implicit device base address when the base itself has not been mapped.

[Ravi Narayanaswamy via Phabricator via Openmp-commits]

RaviNarayanaswamy 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
----------------
ye-luo wrote:
> 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.
When doing look up cannot have same base for 2 different sections of the array  
See
n
17 2. The corresponding pointer variable becomes an attached pointer for the corresponding list item.
18 3. If the original base pointer and the corresponding attached pointer share storage, then the
19 original list item and the corresponding list item must share storag


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D117997