[flang-commits] [mlir] [openmp] [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:

> Thank you for all your effort in trying to make reviewers happy!

It's no problem at all! 

> Yes - it is better - at least from a type perspective. I am not sure I completely understand why var_ptr had to be made optional (the data action uses var_ptr to know which "host pointer" needs a device pointer - the var_ptr_ptr is for that offset determination in parent). Possibly because of how LLVM decomposes the actions. 

In this case it's optional as the lowering is setup to treat them as a pointer/base pointer coupling at the moment (as I'd only really thought about using the var_ptr field in this PR, as we haven't touched on using the var_ptr_ptr field yet in the lowering for other map related changes) which is from what I am now gathering not the original intent, and I have misunderstood it (part of the reason I'd like a base pointer field to work with at the moment, but I am likely just being quite dense)!

I am more than happy to change it to be non-optional in this PR if we desire to (I'd stupidly misinterpreted it as being a one or the other thing from your previous comment), but it may require the addition of a base pointer field (perhaps not though, I do have an idea to deal with that a little better, but I am not sure it'll pan out). 

For the case of the descriptor base address in the example in the document (in which we've moved it from the var_ptr to var_ptr_ptr field), what would then go in the var_ptr field to make it a non-optional component if you do not mind me asking?  

>But alas, I think at this point you probably need to move this forward. So consider an approval from me - I won't formally approve as I feel that requires more due diligence to look into detail in all parts of your change.
> 

It's no problem, I'd like to get it into a state we can all be at least reasonably happy with until the next iteration :-) As a lot of this will be revisited in more detail to get the lowering and map info in a more complete and generalised state for derived types (and other record types from languages like C/C++ hopefully). I do not mind delaying it further if that's what everybody wishes, I'd just need some direction on what to touch on to improve in this case for the next review cycle, I'll leave it up to @kiranchandramohan to decide what he wishes done in this case.

> Regarding the renaming of the fields, I can understand your explanation from the point-of-view of consistency with OpenMPIRBuilder (and the OpenMP spec of base pointer). However, I don't necessarily think anything needs removed. Namely each existing piece is useful as far as I can see and base pointer maybe could be in addition.
> 

That's another possibility one I'd be happy with as well, I am primarily trying to work out if either change is necessary as I work through trying to map derived types! There's a possibility it's not, I just need to work through the various cases (primarily just focusing on partial mapping of derived types via explicit member mapping for the first PR), and the way the fields are currently utilised in the lowering (or at least in my mind) they seemed synonymous with base pointer, and offload pointer. 

> Namely, the existing var_ptr is more precise than base_pointer since it allows partial mapping. Consider the example in the OpenMP standard for data terminology: COMMENT: For the array section (*p0).x0[k1].p1->p2[k2].x1[k3].x2[4][0:n], where identifiers pi have a pointer type declaration and identifiers xi have an array type declaration, the base pointer is: (*p0).x0[k1].p1->p2. var_ptr can capture the address of "(*p0).x0[k1].p1->p2[k2].x1[k3].x2[4]"

I see what you mean in this case, and I believe I do use the var_ptr field for that, as it's more akin to how OpenMP's (or OpenMPIRBuilder/libomptarget, as I don't think there's a definition in the specification for this one, attached pointer may be the closest) offload pointer works I believe. It's just in certain cases (primarily anything to do with a parent <-> child relationship from what I've found so far) we need to also know the base pointer as well to appropriately offload and I don't think it makes sense (to at least me, but perhaps I am wrong, always likely) to have to dig it up in the lowering from the LLVM-IR (seems like it could be quite prone to failure at least, but maybe I am overthinking it) when the original generator of the MLIR should have the necessary information to specify what we need. And I originally mistook var_ptr_ptr to be the equivalent of that but it appears I am rather wrong in this case! :-)  

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