[llvm] [OFFLOAD] Fix an issue with nested close, alloc mapping when using USM (PR #208122)

[Deepak Eachempati via llvm-commits]

dreachem wrote:

> However, there is a case where this optimization is not compatible even with 6.0:
> 
> ```c
>   #pragma omp requires unified_shared_memory
> 
>   int leafarr[10];
>   struct Mid { int *p2; };
>   struct Mid mid;
>   mid.p2 = &leafarr[0];
>   struct Mid *p1 = ∣          // p1 -> mid, and mid.p2 -> leafarr
> 
>   #pragma omp target enter data map(alloc: p1)
>   // device(p1) == host(p1)
>   // printf("%d\n", omp_get_mapped_ptr(&p1, ...) == &p1); // 1
> 
>   #pragma omp target enter data map(close, to: leafarr[0:10])
>   // device(leafarr) != host(leafarr)
>   // printf("%d\n", omp_get_mapped_ptr(&leafarr, ...) == &leafarr); // 0
> 
> 
>   #pragma omp target enter data map(present, alloc: p1)      \
>                                 map(close,   alloc: p1[0:1]) \ // new allocation
>                                 map(present, alloc: p1->p2[0:10])
>   // mid (which CONTAINS p2) is being newly mapped, so we must choose one backing
>   // for it, and both choices result in bad pointer-attachments:
>   //
>   //   * device(p1[0:1]) == host(p1[0:1])
>   //     - OK:  host(p1) => host(p1[0])
>   //     - BAD: host(p1[0].p2) => device(p1[0].p2[0:10]) 
>   //
>   //   * device(p1[0:1]) != host(p1[0:1])
>   //     - BAD: host(p1) => device(p1[0])
>   //     - OK:  host(p1[0].p2) => host(p1[0].p2[0:10]);
> ```
> 
> So in such a scenario, the runtime would have no option but to emit a runtime error, if this optimization is enabled, which it is, by default. So, I don't see a way to make the implementation/optimization fully leagal even for 6.0.

That's an interesting example. Seems like a hole that the spec should address. The two options would be emitting a runtime error or not doing the second pointer attachment (since the pointer attachment "shared pointee" condition cannot be satisfied). 

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