[llvm] [OpenMP][Offload] Handle non-memberof `present/to/from` entries irrespective of order. (PR #165494)

[Abhinav Gaba via llvm-commits]

abhinavgaba wrote:

We don't have to worry about cases like `map(alloc: a[0:4]) map(to: a[1])` since OpenMP guarantees that the two list-items should be identical in this case.

_OpenMP 6.0 page 286:_

```none
18 • Two list items of the map clauses on the same construct must not share original storage
19   unless one of the following is true: they are the same list item, one is the containing structure
20   of the other, at least one is an assumed-size array, or at least one is implicitly mapped due to
21   the list item also appearing in a use_device_addr clause.
```

However, the following is allowed. @dreachem might be able to confirm:

```c
 int *p = ...;
 #pragma omp target map(to : p[0:4]) map(present, from: p[0:4]) map(delete, storage: p[:])
```

Which can be handled by re-ordering for the `enter_data` part similar to (1) , but not the `exit_data` part:
```
// (3.1)
#pragma omp target_enter_data map(present, from: p[0:4]) // "from" decays into "storage"
                              map(to: p[0:4]) // No transfer because "present"
                              map(delete, storage: p[:])) // No-op

// ...
// (3.2)
#pragma omp target_exit_data map(delete, storage: p[:]))
                             map(present, from: p[0:4])
                             map(to: p[0:4]) // "to" decays into "storage"
```
There is no clear way to order "delete" vs "from" if they are on different maps.

This can still be handled by the compiler like the following, but this case makes it more complicated for the compiler since the expressions are not identical.

```
// Possible compiler codegen
// (4.1)
#pragma omp target_enter_data map(present, to: p[0:4]) // use list-item from "to"
                                                       // and add "present".

...
// (4.2)
#pragma omp target_exit_data map(present, delete, from: p[0:4])) // Use list-item from "from"
                                                                 // and add "present", "delete"
```

-----

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.

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