alexey-bataev 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
> ```
Ok for this particular case.
>
> Even for cases where both the pointer, pointee are mapped, the expectation is
> for the pointer to still be captured by-ref.
Why?
> 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 enter data 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, omp_get_default_device());
>
> #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.");
> }
> ```
>
In this particular case we may pass it byref, but not necessary in all cases
> 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
> [abhinavgaba#1](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
