[llvm-branch-commits] [flang] [mlir] [Flang][OpenMP] Add pass to replace allocas with device shared memory (PR #161863)

[Michael Kruse via llvm-branch-commits]

https://github.com/Meinersbur commented:

> Similarly, any random dialect's op lowering to `llvm.alloca` would end up having to be patched too, and doing context-dependent `llvm.alloca` translation to OpenMP device shared memory APIs I think we all agree would be entirely out of the question, since we don't want to add a dependency to OpenMP to the LLVM dialect. There's nothing we could do at the OpenMPToLLVMIRTranslation stage, since that can only deal with `omp` dialect operations and discardable attributes there.

Makes sense. Invoking the pass later before lowering is a good idea.

One question though: How do we know that the semantics of a particular `llvm.alloca` is meant to be shared memory? This PR is inferring it from its uses, but this doesn't necessarily mean oit is the intition. For instace, the PR has exceptions for `private_syms`, but how dow we know there are no other exceptions?


> > To be usable in indirectably reached regions, the adress mlir::Value has to be passed to it in either case, why is it relevant?
> 
> Unfortunately, I don't think I quite understand this concern. Could you elaborate on it?

It's been some time ago, but I think I meant something like this:
```c
void foobar(void *x) {
  // access *x
}

#pragma omp target private(x) {
{
  #pragma omp teams {
    foobar(&x);
  }
}
```
`foobar` (the "indirectably reached region") could not access the shared_alloca `x` without having it passed as a pointer as an argument. And `foobar` does not need to know where the memory comes from. It was referring to the second paragraph of the .td doc of StackToShared.

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