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

[Ye Luo via Phabricator via Openmp-commits]

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
----------------
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.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D117997