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

[Razvan Lupusoru 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 = ""}
----------------
razvanlupusoru wrote:

Why is the var_ptr a pointer to a pointer? This inconsistency feels unintuitive to me. A mapping operation is intended to map the data pointed to by a pointer recorded in `var_ptr`. In this case, it feels it is mapping a pointer instead of the data itself.

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