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

[Deepak Eachempati via llvm-commits]

dreachem wrote:

@abhinavgaba 

> 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:
> 
> ```c
> // (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.

The "delete" should just be handled at the end, after everything else. I'm not too familiar with the LLVM runtime implementation, so it's not clear to me what the technical difficulty is here. Is there a reason the algorithm in the map clause section (based on "mappable storage blocks") can't be used? It does require grouping individual maps into sets that apply to a common storage block. These are the pertinent sentences:

> All map clause list items that share storage or have the same containing structure or containing array result in a single mappable storage block that contains the storage of the list items, unless otherwise specified. The storage for each other map clause list item becomes a distinct mappable storage block. If a list item is a referencing variable that has a containing structure, the behavior is as if only the storage for its referring pointer is part of that structure. In general, if a list item is a referencing variable then the storage for its referring pointer and its referenced pointee occupy distinct mappable storage blocks.

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

> If multiple list items are explicitly mapped on the same construct and have the same containing array or have base pointers that share original storage, and if any of the list items do not have corresponding list items that are present in the device data environment prior to a task encountering the construct, then the list items must refer to the same array elements of either the containing array or the implicit array of the base pointers.

In general, the runtime would need to be involved in some of the matching because we allow overlapping storage for assumed-size arrays and use_device_addr. Other than these cases, the compiler should in theory be able to group the maps into sets where each set applies to the same storage block.

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