[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:
> RaviNarayanaswamy wrote:
> > 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?
> Ravi, I was thinking of
> ```
> #include <iostream>
> 
> int main()
> {
>   int a[30];
>   a[ 1] = 0;
>   a[21] = 0;
>   #pragma omp target map(from: a[:10], a[20:10])
>   {
>     a[ 1] = 1;
>     a[21] = 2;
>   }
>   std::cout << a[1] << " " << a[21] << std::endl;
> }
> ```
> But probably this code is illegal.
  #pragma omp target map(from: a[:10], a[20:10])  is illegal
 


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D117997