[llvm] [OpenMP][Offload] Handle non-memberof `present/to/from` entries irrespective of order. (PR #165494)
Deepak Eachempati via llvm-commits
llvm-commits at lists.llvm.org
Fri Oct 31 09:06:34 PDT 2025
dreachem wrote:
@abhinavgaba
> Another thing to think of, but maybe from the spec's perspective, is that should the modifiers from a `from/to` clause be completely ignored on `enter/exit_data`?
>
> For the following:
>
> ```c
> int x;
>
> int main(void) {
> #pragma omp target map(delete, to: x) map(present, from: x)
> ;
> }
> ```
>
> @dreachem, @mjklemm, is this expected to run OK? If we decay `from` into `storage`, and then apply it to `enter_data`, then it would seem to imply that it's equivalent to doing `map(present, to: x)`, in which case, this will cause a runtime-error-termination.
>
> ```c
> // Applying "present" to "enter_data" would cause runtime-error
> // (5.1)
> #pragma omp target_enter_data(present, to: x)
> ...
> // (5.2)
> #pragma omp target_exit_data(present, delete, from: x)
> ```
>
> But the user may be expecting to only do a present-check on `exit_data` here, which would means we need to ignore any modifiers from a `from/to` clause that's decayed into `storage` on `target_enter/exit_data`.
>
> ```c
> // Ignoring modifiers from the decayed-to-storage clause.
> // (6.1)
> #pragma omp target_enter_data(to: x) // delete ignored for enter_data
> ...
> // (6.2)
> #pragma omp target_exit_data(present, from: x)
> ```
>
> But if we do so, then `delete` gets ignored on `target_exit_data`, since it came from a decayed `to`. Maybe we only special-case present and not apply it to decayed-to-storage maps.
So, you are describing `target` as if it there is a `target_enter_data` on entry and a `target_exit_data` on exit, but this isn't how the specification describes it. That's the description for `target_data`.
For `target`, a `map(delete, to: x)` clause has the effect of `map(to: x)` on entry and `map(delete,storage:x)` on exit, and a `map(present,from:x)` clause has the effect of a `map(present,storage:x)` on entry and a `map(from:x)` on exit. So, putting it together, for that `target` construct you effectively get the equivalent of `map(present, to: x)` on entry and `map(delete,from:x)` on exit.
Let's suppose the construct was instead `target_data map(delete,to: x) map(present,from: x)`. In that case, we describe it as equivalent to:
```
// note that "delete" on target_enter_data is just ignored
#pragma omp target_enter_data map(delete,to: x) map(present,storage: x)
#pragma omp task if(0) // a merged "sharing task"
;
#pragma omp target_exit_data map(delete,storage: x) map(present,from: x)
```
Here, it looks like we hit a possible unintended discrepancy. When we decided to make `target_data` composite, I'm not sure we accounted for the behavior of the `present` modifier possibly changing. Since there are now 2 data-mapping regions each with their own entry and exit, as opposed to a single `target_data` region with one entry and one exit, there are two points at which the `present` modifier can have an effect. I don't think this behavior is unreasonable for `target_data`, but just wanted to call it out.
https://github.com/llvm/llvm-project/pull/165494
More information about the llvm-commits
mailing list