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

[Abhinav Gaba via llvm-commits]

abhinavgaba wrote:

I found multiple issues. I created a draft PR #213214, that includes this branch, with some follow-up changes and tests, trying to try some fixes.

The main issue is that there is information like the size of `declare_target` globals, that the runtime simply doesn't get under USM, and without that, there's no way for the runtime to know whether a variable being mapped with `close` falls within the storage block of something that was marked as `declare_target`. So that requires codegen change on the compiler side.


-----

The codegen change in that PR fixes `declare_target` issues. However, the issue I'm struggling with is about pointer-attachment. I think that needs clarification from OpenMP spec.

There can be cases like:
```c
#pragma omp requires unified_shared_memory
...
int x[10];
int *p = &x[0];
#pragma omp target_enter_data map(alloc: p) // device(p) = host(p)

#pragma omp target_enter_data map(close, alloc: p[0:10]) // device(p[0:10]) != host(p[0:10])
// Attaches device(p) => device(p[0:10])
// But device(p) == host(p) as we didn't allocate a separate "device(p)"
// So, this does host(p) => device(p[0:10])

printf("%d\n", p[0]); // Segfault: dereferences device(p[0]) on host
...
```

Or even worse, where pointee and pointer are both already allocated, and then attached later on:

```c
int x[10];
int *p = &x[0];
#pragma omp target_enter_data map(alloc: p) // device(p) = host(p)
#pragma omp target_enter_data map(close, alloc: x[0:10]) // device(p[0:10]) != host(p[0:10])

#pragma omp target_enter_data map(attach(always): p[0:10]) // device(p) => device(p[0:10])
// But device(p) == host(p) as we didn't allocate a separate "device(p)"
// So, this does host(p) => device(p[0:10])

printf("%d\n", p[0]); // Segfault: dereferences device(p[0]) on host
...

```

>From OpenMP's point of view, we say that "corresponding list-item" may share storage with the "original list-item", so it's not clear to me whether this dereference through the host-pointer `p`, while its corresponding device pointer is in an "attached" state, is expected to "work", or if this is implementation-defined behavior, or it's a user-error. @dreachem might have some insight on this.

If it's supposed to work, then this clang optimization to use actual device memory for `map(close, *: x)` and use host memory as corresponding-storage for everything else, largely becomes illegal, as it means anything that's a pointer, or contains a pointer, always needs to have its own corresponding device storage that's different from original/host-storage, because eventually, some device-pointee may attach to it.

```c
struct S {
  int x;
  int y;
  int *p;
};

S s;
char *buf = &s;

#pragma omp target_enter_data map(alloc: buf[0:sizeof(S)]) // device(s) == host(s)

#pragma omp target_enter_data map(close, alloc: s.p[0:10]) // device(s.p[0:10]) != host(s.p[0:10])

#pragma omp target_enter_data map(attach(always): s.p[0:10])
// device(s.p) => device(s.p[0:10])
// === host(s.p) => device(s.p[0:10])

printf("%d\n", s.p[0]); // Segfault: dereferences device(p[0]) on host
```

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