[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