This is an archive of the discontinued LLVM Phabricator instance.

[libomptarget] Correctly return the implicit device base address when the base itself has not been mapped.
Needs ReviewPublic

Authored by grokos on Jan 23 2022, 10:44 AM.

Details

Summary

When an array is mapped starting from an offset, its base address does not fall within the mapped range. If the same array is later captured implicitly in a target region or we use a use_device_ptr clause, the array appears as unmapped. This leads to otherwise legal OpenMP programs failing to execute correctly.

Diff Detail

Event Timeline

grokos requested review of this revision.Jan 23 2022, 10:44 AM
grokos created this revision.
ye-luo added inline comments.Jan 23 2022, 11:06 AM
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

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.

grokos marked an inline comment as done.Jan 23 2022, 12:45 PM
grokos added inline comments.
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

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.

ABataev added inline comments.Jan 23 2022, 12:50 PM
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

IIRC, Ye is correct here. Mapping 2 sections in different constructs is legal, just the inner one must be a subset of the outer.

grokos marked an inline comment as done.Jan 23 2022, 1:07 PM
grokos added inline comments.
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

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.

ye-luo added inline comments.Jan 23 2022, 2:02 PM
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

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.

grokos updated this revision to Diff 402547.Jan 24 2022, 8:38 AM
grokos retitled this revision from [libomptarget] Lookup by base address when begin address is not mapped. to [libomptarget] Correctly return the implicit device base address when the base itself has not been mapped..

Added test for use_device_ptr

Changed the title because it was totally misleading.

RaviNarayanaswamy added inline comments.
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

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

openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

Ye, are you asking if there are 2 sections mapped form same array ie they have same base then the look up should fail?

ye-luo added inline comments.Jan 24 2022, 3:00 PM
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32

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.

openmp/libomptarget/test/mapping/array_section_implicit_capture.c
32
#pragma omp target map(from: a[:10], a[20:10])  is illegal
grokos updated this revision to Diff 403626.Jan 27 2022, 6:09 AM

Added a check to catch the error when the user tries to map distinct sections of the same object using different pointers.