[clang] [llvm] [Clang][OpenMP] Capture mapped pointers on `target` by reference. (PR #145454)

[Abhinav Gaba via llvm-commits]

abhinavgaba wrote:

> I have one big question - why do we need this? If pointer is mapped explicitly, it becomes an attached pointer, which cannot be modified in target region. How we capture/process it, byref or byval, does not matter, we can do whatever is better. So, why do we need to process it byref?

That's a good question. There are multiple reasons.

First, just mapping a pointer does not make it attached. For a pointer to become attached, a map has to be on the pointee, like `map(p[0:1])`, and then the attachment should only happen when either the pointer is new, or the pointee is new on that construct on which the pointee is being mapped. And whether an attachment will be successful or not cannot be determined at compile-time (unless the user uses OpenMP 6.1's `attach(always)`).

So, a user must be able to access and modify a pointer in a target region, and bring it back, just like any variable, as long as it is not attached.

```c
int *p;
#pragma omp target map(tofrom:p) // p is not an attached
  p = (int*) 111;

printf("%d\n", (int) p); // CHECK: 111
```

Even for cases where both the pointer, pointee are mapped, the expectation is for the pointer to still be captured by-ref. If the attachment is not supposed to happen (based on the conditions mentioned above), then any modifications to the pointer have to be brought back if the map-type is from.

```c
int *p = ...;
#pragma omp target map(tofrom:p[0:1]) // p is not attached, since p is not present on the device
#pragma omp target map(tofrom:p) // p is not attached, since pointee is not mapped on this construct
  p = (int*) 111;

printf("%d\n", (int) p); // CHECK: 111 
```

And irrespective of whether the attachment succeeds or not, a user can assert that the address of the pointer in the target region has to match its mapped address. e.g.:

```c
int *p;
#pragma omp target enter data map(p)
int64_t mapped_p = (int64_t) omp_get_mapped_ptr(&p);

#pragma omp target map(present,alloc: p) map(p[0:10])
{
 assert(mapped_p == (int64_t) &p && "address of p in device should match the mapped address before the region.");
}
```

Additionally, before this change, when we are using PTR_AND_OBJ when both pointer/pointee are mapped (which is not correct as it does not follow the conditional pointer-attachment), we are capturing the pointer by value. With the ATTACH mapping support in https://github.com/abhinavgaba/llvm-project/pull/1, we aren't using PTR_AND_OBJ anymore when the base-pointer is a non-struct-member pointer variable. We instead map the pointer and pointee with independent maps and then attach them conditionally. With that, the kernel signature needs to match for `map(p)` and `map(p, p[0:1])`, because for both of them, the PARAM should correspond to the same `map(p)`.

https://github.com/llvm/llvm-project/pull/145454