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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
