[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 openmp-commits at lists.llvm.org
Mon Jan 24 14:49:25 PST 2022

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
RaviNarayanaswamy wrote:
> 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
Ye,  are you  asking if there are 2 sections mapped form same array ie they have same base then the look up should fail?

  rG LLVM Github Monorepo



More information about the Openmp-commits mailing list