[flang-commits] [openmp] [mlir] [flang] [Flang][OpenMP] Initial mapping of Fortran pointers and allocatables for target devices (PR #71766)

[via flang-commits]

================
@@ -0,0 +1,125 @@
+<!--===- docs/OpenMP-descriptor-management.md
+
+   Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+   See https://llvm.org/LICENSE.txt for license information.
+   SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+-->
+
+# OpenMP dialect: Fortran descriptor type mapping for offload
+
+The initial method for mapping Fortran types tied to descriptors for OpenMP offloading is to treat these types 
+as a special case of OpenMP record type (C/C++ structure/class, Fortran derived type etc.) mapping as far as the 
+runtime is concerned. Where the box (descriptor information) is the holding container and the underlying 
+data pointer is contained within the container, and we must generate explicit maps for both the pointer member and
+the container. As an example, a small C++ program that is equivalent to the concept described, with the 
+`mock_descriptor` class being representative of the class utilised for descriptors in Clang:
+
+```C++
+struct mock_descriptor {
+  long int x;
+  std::byte x1, x2, x3, x4;
+  void *pointer;
+  long int lx[1][3];
+};
+
+int main() {
+mock_descriptor data;
+#pragma omp target map(tofrom: data, data.pointer[:upper_bound])
+{
+    do something... 
+}
+
+ return 0;
+}
+```
+
+In the above, we have to map both the containing structure, with its non-pointer members and the
+data pointed to by the pointer contained within the structure to appropriately access the data. This 
+is effectively what is done with descriptor types for the time being. Other pointers that are part 
+of the descriptor container such as the addendum should also be treated as the data pointer is 
+treated.
+
+Currently, Flang will lower these descriptor types in the OpenMP lowering (lower/OpenMP.cpp) similarly
+to all other map types, generating an omp.MapInfoOp containing relevant information required for lowering
+the OpenMP dialect to LLVM-IR during the final stages of the MLIR lowering. However, after 
+the lowering to FIR/HLFIR has been performed an OpenMP dialect specific pass for Fortran, 
+`OMPDescriptorMapInfoGenPass` (Optimizer/OMPDescriptorMapInfoGen.cpp) will expand the 
+`omp.MapInfoOp`'s containing descriptors (which currently will be a `BoxType` or `BoxAddrOp`) into multiple 
+mappings, with one extra per pointer member in the descriptor that is supported on top of the original
+descriptor map operation. These pointers members are linked to the parent descriptor by adding them to 
+the member field of the original descriptor map operation, they are then inserted into the relevant map
+owning operation's (`omp.TargetOp`, `omp.DataOp` etc.) map operand list and in cases where the owning operation
+is `IsolatedFromAbove`, it also inserts them as `BlockArgs` to canonicalize the mappings and simplify lowering.
+
+An example transformation by the `OMPDescriptorMapInfoGenPass`:
+
+```
+
+...
+%12 = omp.map_info var_ptr(%1#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.box<!fir.ptr<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%11) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> {name = "arg_alloc"}
+...
+omp.target map_entries(%12 -> %arg1, %13 -> %arg2 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<i32>) {
+    ^bb0(%arg1: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, %arg2: !fir.ref<i32>):
+...
+
+====>
+
+...
+%12 = fir.box_offset %1#1 base_addr : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+%13 = omp.map_info var_ptr(%12 : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%11) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
----------------
agozillon wrote:

> > (although, it could be modified to utilise var_ptr_ptr when it's provided in place of var_ptr, but that seems like a band-aid).
> 
> I am not sure I understand - it doesn't seem like a band-aid to me. The current approach seems like an issue to me where mapping semantics are overloaded and the semantics no longer mean map the data pointed to.

In the sense that I'd like to change the semantics of var_ptr/var_ptr_ptr -> base pointer and pointer, it means that I'd be modifying it again in the near future. But it's not a problem in any case.

> > I am hoping to change that soon in a future PR, but I am also considering changing the name to base pointer/offload pointer (or perhaps base pointer would be better as a new field)
> 
> I am curious what you have in mind for this and whether it is more intuitive than the current model. And just to be clear, I prefer convincing in the form that the new IR is easier to understand/optimize, not that it matches LLVM's OpenMPIRBuilder. IMO, the omp dialect should not have coupling with the IR builder that would make it harder to optimize and lower through MLIR's GPU pipeline.
> 

It's easier to understand from the perspective of it is what they're being used as currently when lowering to LLVM-IR, alongside it's also the terminology used by the OpenMP specification, Clang and libomptarget. So I would imagine it'd be a lot clearer to readers and implementers alike what the intent is. In this case I do not see how it'd make it any easier or harder to work with MLIR's GPU pipeline, especially as the current suggested combination of the GPU+OMP dialect seems to utilise the OMPIRBuilder. However, if you think it's clearer as is, that's fine, we can see as we progress.

> > No worries, either fixing it in this PR or delaying it works for me, but as I said.
> 
> I am OK with whatever you choose. The initial fix seems straight forward - just use var_ptr_ptr instead of var_ptr.

Hopefully that should be the case. 


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