[flang-commits] [flang] [flang][NFC] Strip trailing whitespace from markdown files (2 of 2) (PR #173253)
Tarun Prabhu via flang-commits
flang-commits at lists.llvm.org
Tue Dec 23 04:09:20 PST 2025
https://github.com/tarunprabhu updated https://github.com/llvm/llvm-project/pull/173253
>From 8a329a441b04410421a79e2e1452b9174d5247a0 Mon Sep 17 00:00:00 2001
From: Tarun Prabhu <tarun at lanl.gov>
Date: Mon, 22 Dec 2025 05:32:26 -0700
Subject: [PATCH] [flang][NFC] Strip trailing whitespace from markdown files (2
of 2)
Strip trailing whitespace from the remaining markdown files. This completes the
sequence of NFC commits performing the cleanup of trailing whitespace from
markdown files.
---
flang/docs/OpenACC-descriptor-management.md | 18 +-
flang/docs/OpenACC.md | 10 +-
flang/docs/OpenMP-declare-target.md | 226 +++++++++---------
flang/docs/OpenMP-descriptor-management.md | 68 +++---
flang/docs/OpenMP-semantics.md | 6 +-
flang/docs/OptionComparison.md | 20 +-
flang/docs/Overview.md | 20 +-
.../docs/ParallelMultiImageFortranRuntime.md | 9 +-
flang/docs/ParameterizedDerivedTypes.md | 2 +-
flang/docs/ParserCombinators.md | 6 +-
flang/docs/Parsing.md | 8 +-
flang/docs/Preprocessing.md | 6 +-
flang/docs/PullRequestChecklist.md | 2 +-
flang/docs/RuntimeDescriptor.md | 6 +-
flang/docs/RuntimeEnvironment.md | 6 +-
flang/docs/RuntimeTypeInfo.md | 6 +-
flang/docs/Semantics.md | 6 +-
flang/docs/Unsigned.md | 6 +-
18 files changed, 215 insertions(+), 216 deletions(-)
diff --git a/flang/docs/OpenACC-descriptor-management.md b/flang/docs/OpenACC-descriptor-management.md
index 7a6e6b4346348..a9618667af387 100644
--- a/flang/docs/OpenACC-descriptor-management.md
+++ b/flang/docs/OpenACC-descriptor-management.md
@@ -53,7 +53,7 @@ program main
real, pointer :: t1(:,:)
nullify(d%p)
allocate(t1(2,2))
-
+
! 2.7.9:
! Allocates the memory for object 'd' on the device.
! The descriptor member of 'd' is a NULL descriptor,
@@ -62,9 +62,9 @@ program main
! The descriptor storage is created on the device
! just as part of the object 'd' storage.
!$acc enter data create(d)
-
+
d%p => t1
-
+
! 2.7.7:
! Pointer d%p is not present on the device, so copyin
! action is performed for the data pointed to by the pointer:
@@ -80,7 +80,7 @@ program main
! from the host values of the corresponding members.
! * The attachment counter of 'd%p' is set to 1.
!$acc enter data copyin(d%p)
-
+
! 2.7.7:
! Pointer d%p is already present on the device, so no copyin
! action is performed.
@@ -89,12 +89,12 @@ program main
! during the previous attachment, only its attachment counter
! is incremented to 2.
!$acc enter data copyin(d%p)
-
+
! 3.2.29:
! The detach action is performed. According to 2.7.2 the attachment
! counter of d%p is decremented to 1.
call acc_detach(d%p)
-
+
! 3.2.29:
! The detach action is performed. According to 2.7.2 the attachment
! counter of d%p is decremented to 0, which initiates an update
@@ -102,7 +102,7 @@ program main
! pointer in the local memory.
! We will discuss this in more detail below.
call acc_detach(d%p)
-
+
! The following construct will fail, because the 'd%p' descriptor's
! base_addr is now the host address not accessible on the device.
! Without the second 'acc_detach' it will work correctly.
@@ -111,7 +111,7 @@ program main
!$acc end serial
```
-Let's discuss in more detail what happens during the second `acc_detach`.
+Let's discuss in more detail what happens during the second `acc_detach`.
OpenACC 2.6.4:
@@ -255,7 +255,7 @@ Due to `d%p` reference in the `present` clause of the `serial` region, the compi
In the case of POINTER dummy argument, if the descriptor storage is not explicitly created in the user code, the pointer attachment may not happen due to 2.7.2:
> 1693 If the pointer var is in shared memory or is not present in the current device memory, or if the
-> 1694 address to which var points is not present in the current device memory, no action is taken.
+> 1694 address to which var points is not present in the current device memory, no action is taken.
Example:
diff --git a/flang/docs/OpenACC.md b/flang/docs/OpenACC.md
index 6d384f93ae216..f1fe69e57bf37 100644
--- a/flang/docs/OpenACC.md
+++ b/flang/docs/OpenACC.md
@@ -1,9 +1,9 @@
-<!--===- docs/Extensions.md
-
+<!--===- docs/Extensions.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
-
+
-->
# OpenACC in Flang
@@ -23,9 +23,9 @@ local:
warning instead of an error as other compiler accepts it.
* The `if` clause accepts scalar integer expression in addition to scalar
logical expression.
-* `!$acc routine` directive can be placed at the top level.
+* `!$acc routine` directive can be placed at the top level.
* `!$acc cache` directive accepts scalar variable.
-* The `!$acc declare` directive accepts assumed size array arguments for
+* The `!$acc declare` directive accepts assumed size array arguments for
`deviceptr` and `present` clauses.
* The OpenACC specification disallows a variable appearing multiple times in
clauses of `!$acc declare` directives for a function, subroutine, program,
diff --git a/flang/docs/OpenMP-declare-target.md b/flang/docs/OpenMP-declare-target.md
index 478db2d004c87..ac7fec5e83495 100644
--- a/flang/docs/OpenMP-declare-target.md
+++ b/flang/docs/OpenMP-declare-target.md
@@ -18,9 +18,9 @@ local:
## Introduction to Declare Target
In OpenMP `declare target` is a directive that can be applied to a function or
-variable (primarily global) to notate to the compiler that it should be
-generated in a particular device's environment. In essence whether something
-should be emitted for host or device, or both. An example of its usage for
+variable (primarily global) to notate to the compiler that it should be
+generated in a particular device's environment. In essence whether something
+should be emitted for host or device, or both. An example of its usage for
both data and functions can be seen below.
```Fortran
@@ -37,9 +37,9 @@ program main
end program
```
-In the above example, we create a variable in a separate module, mark it
-as `declare target` and then map it, embedding it into the device IR and
-assigning to it.
+In the above example, we create a variable in a separate module, mark it
+as `declare target` and then map it, embedding it into the device IR and
+assigning to it.
```Fortran
@@ -57,28 +57,28 @@ end program
```
In the above example, we are stating that a function is required on device
-utilising `declare target`, and that we will not be utilising it on host,
-so we are in theory free to remove or ignore it there. A user could also
-in this case, leave off the `declare target` from the function and it
-would be implicitly marked `declare target any` (for both host and device),
+utilising `declare target`, and that we will not be utilising it on host,
+so we are in theory free to remove or ignore it there. A user could also
+in this case, leave off the `declare target` from the function and it
+would be implicitly marked `declare target any` (for both host and device),
as it's been utilised within a target region.
## Declare Target as represented in the OpenMP Dialect
-In the OpenMP Dialect `declare target` is not represented by a specific
-`operation`. Instead, it's an OpenMP dialect specific `attribute` that can be
-applied to any operation in any dialect, which helps to simplify the
-utilisation of it. Rather than replacing or modifying existing global or
+In the OpenMP Dialect `declare target` is not represented by a specific
+`operation`. Instead, it's an OpenMP dialect specific `attribute` that can be
+applied to any operation in any dialect, which helps to simplify the
+utilisation of it. Rather than replacing or modifying existing global or
function `operations` in a dialect, it applies to it as extra metadata that
-the lowering can use in different ways as is necessary.
+the lowering can use in different ways as is necessary.
-The `attribute` is composed of multiple fields representing the clauses you
-would find on the `declare target` directive i.e. device type (`nohost`,
-`any`, `host`) or the capture clause (`link` or `to`). A small example of
+The `attribute` is composed of multiple fields representing the clauses you
+would find on the `declare target` directive i.e. device type (`nohost`,
+`any`, `host`) or the capture clause (`link` or `to`). A small example of
`declare target` applied to a Fortran `real` can be found below:
```
-fir.global internal @_QFEi {omp.declare_target =
+fir.global internal @_QFEi {omp.declare_target =
#omp.declaretarget<device_type = (any), capture_clause = (to)>} : f32 {
%0 = fir.undefined f32
fir.has_value %0 : f32
@@ -87,12 +87,12 @@ fir.global internal @_QFEi {omp.declare_target =
This would look similar for function style `operations`.
-The application and access of this attribute is aided by an OpenMP Dialect
-MLIR Interface named `DeclareTargetInterface`, which can be utilised on
+The application and access of this attribute is aided by an OpenMP Dialect
+MLIR Interface named `DeclareTargetInterface`, which can be utilised on
operations to access the appropriate interface functions, e.g.:
```C++
-auto declareTargetGlobal =
+auto declareTargetGlobal =
llvm::dyn_cast<mlir::omp::DeclareTargetInterface>(Op.getOperation());
declareTargetGlobal.isDeclareTarget();
```
@@ -100,127 +100,127 @@ declareTargetGlobal.isDeclareTarget();
## Declare Target Fortran OpenMP Lowering
The initial lowering of `declare target` to MLIR for both use-cases is done
-inside of the usual OpenMP lowering in flang/lib/Lower/OpenMP.cpp. However,
-some direct calls to `declare target` related functions from Flang's
+inside of the usual OpenMP lowering in flang/lib/Lower/OpenMP.cpp. However,
+some direct calls to `declare target` related functions from Flang's
lowering bridge in flang/lib/Lower/Bridge.cpp are made.
-The marking of operations with the declare target attribute happens in two
-phases, the second one optional and contingent on the first failing. The
-initial phase happens when the declare target directive and its clauses
+The marking of operations with the declare target attribute happens in two
+phases, the second one optional and contingent on the first failing. The
+initial phase happens when the declare target directive and its clauses
are initially processed, with the primary data gathering for the directive and
clause happening in a function called `getDeclareTargetInfo`. This is then used
-to feed the `markDeclareTarget` function, which does the actual marking
+to feed the `markDeclareTarget` function, which does the actual marking
utilising the `DeclareTargetInterface`. If it encounters a variable or function
that has been marked twice over multiple directives with two differing device
types (e.g. `host`, `nohost`), then it will swap the device type to `any`.
-Whenever we invoke `genFIR` on an `OpenMPDeclarativeConstruct` from the
-lowering bridge, we are also invoking another function called
+Whenever we invoke `genFIR` on an `OpenMPDeclarativeConstruct` from the
+lowering bridge, we are also invoking another function called
`gatherOpenMPDeferredDeclareTargets`, which gathers information relevant to the
-application of the `declare target` attribute. This information
-includes the symbol that it should be applied to, device type clause,
+application of the `declare target` attribute. This information
+includes the symbol that it should be applied to, device type clause,
and capture clause, and it is stored in a vector that is part of the lowering
-bridge's instantiation of the `AbstractConverter`. It is only stored if we
-encounter a function or variable symbol that does not have an operation
-instantiated for it yet. This cannot happen as part of the
-initial marking as we must store this data in the lowering bridge and we
-only have access to the abstract version of the converter via the OpenMP
-lowering.
-
-The information produced by the first phase is used in the second phase,
-which is a form of deferred processing of the `declare target` marked
-operations that have delayed generation and cannot be proccessed in the
-first phase. The main notable case this occurs currently is when a
-Fortran function interface has been marked. This is
-done via the function
+bridge's instantiation of the `AbstractConverter`. It is only stored if we
+encounter a function or variable symbol that does not have an operation
+instantiated for it yet. This cannot happen as part of the
+initial marking as we must store this data in the lowering bridge and we
+only have access to the abstract version of the converter via the OpenMP
+lowering.
+
+The information produced by the first phase is used in the second phase,
+which is a form of deferred processing of the `declare target` marked
+operations that have delayed generation and cannot be proccessed in the
+first phase. The main notable case this occurs currently is when a
+Fortran function interface has been marked. This is
+done via the function
`markOpenMPDeferredDeclareTargetFunctions`, which is called from the lowering
-bridge at the end of the lowering process allowing us to mark those where
-possible. It iterates over the data previously gathered by
-`gatherOpenMPDeferredDeclareTargets`
-checking if any of the recorded symbols have now had their corresponding
-operations instantiated and applying the declare target attribute where
-possible utilising `markDeclareTarget`. However, it must be noted that it
-is still possible for operations not to be generated for certain symbols,
-in particular the case of function interfaces that are not directly used
-or defined within the current module. This means we cannot emit errors in
-the case of left-over unmarked symbols. These must (and should) be caught
+bridge at the end of the lowering process allowing us to mark those where
+possible. It iterates over the data previously gathered by
+`gatherOpenMPDeferredDeclareTargets`
+checking if any of the recorded symbols have now had their corresponding
+operations instantiated and applying the declare target attribute where
+possible utilising `markDeclareTarget`. However, it must be noted that it
+is still possible for operations not to be generated for certain symbols,
+in particular the case of function interfaces that are not directly used
+or defined within the current module. This means we cannot emit errors in
+the case of left-over unmarked symbols. These must (and should) be caught
by the initial semantic analysis.
NOTE: `declare target` can be applied to implicit `SAVE` attributed variables.
However, by default Flang does not represent these as `GlobalOp`'s, which means
we cannot tag and lower them as `declare target` normally. Instead, similarly
-to the way `threadprivate` handles these cases, we raise and initialize the
+to the way `threadprivate` handles these cases, we raise and initialize the
variable as an internal `GlobalOp` and apply the attribute. This occurs in the
flang/lib/Lower/OpenMP.cpp function `genDeclareTargetIntGlobal`.
## Declare Target Transformation Passes for Flang
-There are currently two passes within Flang that are related to the processing
+There are currently two passes within Flang that are related to the processing
of `declare target`:
* `MarkDeclareTarget` - This pass is in charge of marking functions captured
(called from) in `target` regions or other `declare target` marked functions as
-`declare target`. It does so recursively, i.e. nested calls will also be
-implicitly marked. It currently will try to mark things as conservatively as
+`declare target`. It does so recursively, i.e. nested calls will also be
+implicitly marked. It currently will try to mark things as conservatively as
possible, e.g. if captured in a `target` region it will apply `nohost`, unless
-it encounters a `host` `declare target` in which case it will apply the `any`
-device type. Functions are handled similarly, except we utilise the parent's
+it encounters a `host` `declare target` in which case it will apply the `any`
+device type. Functions are handled similarly, except we utilise the parent's
device type where possible.
* `FunctionFiltering` - This is executed after the `MarkDeclareTarget`
pass, and its job is to conservatively remove host functions from
-the module where possible when compiling for the device. This helps make
-sure that most incompatible code for the host is not lowered for the
-device. Host functions with `target` regions in them need to be preserved
-(e.g. for lowering the `target region`(s) inside). Otherwise, it removes
-any function marked as a `declare target host` function and any uses will be
-replaced with `undef`'s so that the remaining host code doesn't become broken.
-Host functions with `target` regions are marked with a `declare target host`
+the module where possible when compiling for the device. This helps make
+sure that most incompatible code for the host is not lowered for the
+device. Host functions with `target` regions in them need to be preserved
+(e.g. for lowering the `target region`(s) inside). Otherwise, it removes
+any function marked as a `declare target host` function and any uses will be
+replaced with `undef`'s so that the remaining host code doesn't become broken.
+Host functions with `target` regions are marked with a `declare target host`
attribute so they will be removed after outlining the target regions contained
inside.
-While this infrastructure could be generally applicable to more than just Flang,
-it is only utilised in the Flang frontend, so it resides there rather than in
-the OpenMP dialect codebase.
+While this infrastructure could be generally applicable to more than just Flang,
+it is only utilised in the Flang frontend, so it resides there rather than in
+the OpenMP dialect codebase.
## Declare Target OpenMP Dialect To LLVM-IR Lowering
-The OpenMP dialect lowering of `declare target` is done through the
-`amendOperation` flow, as it's not an `operation` but rather an
+The OpenMP dialect lowering of `declare target` is done through the
+`amendOperation` flow, as it's not an `operation` but rather an
`attribute`. This is triggered immediately after the corresponding
operation has been lowered to LLVM-IR. As it is applicable to
-different types of operations, we must specialise this function for
-each operation type that we may encounter. Currently, this is
+different types of operations, we must specialise this function for
+each operation type that we may encounter. Currently, this is
`GlobalOp`'s and `FuncOp`'s.
-`FuncOp` processing is fairly simple. When compiling for the device,
-`host` marked functions are removed, including those that could not
-be removed earlier due to having `target` directives within. This
-leaves `any`, `device` or indeterminable functions left in the
-module to lower further. When compiling for the host, no filtering is
-done because `nohost` functions must be available as a fallback
+`FuncOp` processing is fairly simple. When compiling for the device,
+`host` marked functions are removed, including those that could not
+be removed earlier due to having `target` directives within. This
+leaves `any`, `device` or indeterminable functions left in the
+module to lower further. When compiling for the host, no filtering is
+done because `nohost` functions must be available as a fallback
implementation.
-For `GlobalOp`'s, the processing is a little more complex. We
-currently leverage the `registerTargetGlobalVariable` and
-`getAddrOfDeclareTargetVar` `OMPIRBuilder` functions shared with Clang.
-These two functions invoke each other depending on the clauses and options
+For `GlobalOp`'s, the processing is a little more complex. We
+currently leverage the `registerTargetGlobalVariable` and
+`getAddrOfDeclareTargetVar` `OMPIRBuilder` functions shared with Clang.
+These two functions invoke each other depending on the clauses and options
provided to the `OMPIRBuilder` (in particular, unified shared memory). Their
-main purposes are the generation of a new global device pointer with a
-"ref_" prefix on the device and enqueuing metadata generation by the
-`OMPIRBuilder` to be produced at module finalization time. This is done
-for both host and device and it links the newly generated device global
+main purposes are the generation of a new global device pointer with a
+"ref_" prefix on the device and enqueuing metadata generation by the
+`OMPIRBuilder` to be produced at module finalization time. This is done
+for both host and device and it links the newly generated device global
pointer and the host pointer together across the two modules.
Similarly to other metadata (e.g. for `TargetOp`) that is shared across
-both host and device modules, processing of `GlobalOp`'s in the device
-needs access to the previously generated host IR file, which is done
-through another `attribute` applied to the `ModuleOp` by the compiler
-frontend. The file is loaded in and consumed by the `OMPIRBuilder` to
-populate it's `OffloadInfoManager` data structures, keeping host and
+both host and device modules, processing of `GlobalOp`'s in the device
+needs access to the previously generated host IR file, which is done
+through another `attribute` applied to the `ModuleOp` by the compiler
+frontend. The file is loaded in and consumed by the `OMPIRBuilder` to
+populate it's `OffloadInfoManager` data structures, keeping host and
device appropriately synchronised.
The second (and more important to remember) is that as we effectively replace
the original LLVM-IR generated for the `declare target` marked `GlobalOp` we
-have some corrections we need to do for `TargetOp`'s (or other region
+have some corrections we need to do for `TargetOp`'s (or other region
operations that use them directly) which still refer to the original lowered
global operation. This is done via `handleDeclareTargetMapVar` which is invoked
as the final function and alteration to the lowered `target` region, it's only
@@ -228,38 +228,38 @@ invoked for device as it's only required in the case where we have emitted the
"ref" pointer , and it effectively replaces all uses of the originally lowered
global symbol, with our new global ref pointer's symbol. Currently we do not
remove or delete the old symbol, this is due to the fact that the same symbol
-can be utilised across multiple target regions, if we remove it, we risk
-breaking lowerings of target regions that will be processed at a later time.
-To appropriately delete these no longer necessary symbols we would need a
-deferred removal process at the end of the module, which is currently not in
-place. It may be possible to store this information in the OMPIRBuilder and
-then perform this cleanup process on finalization, but this is open for
+can be utilised across multiple target regions, if we remove it, we risk
+breaking lowerings of target regions that will be processed at a later time.
+To appropriately delete these no longer necessary symbols we would need a
+deferred removal process at the end of the module, which is currently not in
+place. It may be possible to store this information in the OMPIRBuilder and
+then perform this cleanup process on finalization, but this is open for
discussion and implementation still.
## Current Support
For the moment, `declare target` should work for:
-* Marking functions/subroutines and function/subroutine interfaces for
+* Marking functions/subroutines and function/subroutine interfaces for
generation on host, device or both.
-* Implicit function/subroutine capture for calls emitted in a `target` region
-or explicitly marked `declare target` function/subroutine. Note: Calls made
-via arguments passed to other functions must still be themselves marked
-`declare target`, e.g. passing a `C` function pointer and invoking it, then
-the interface and the `C` function in the other module must be marked
-`declare target`, with the same type of marking as indicated by the
+* Implicit function/subroutine capture for calls emitted in a `target` region
+or explicitly marked `declare target` function/subroutine. Note: Calls made
+via arguments passed to other functions must still be themselves marked
+`declare target`, e.g. passing a `C` function pointer and invoking it, then
+the interface and the `C` function in the other module must be marked
+`declare target`, with the same type of marking as indicated by the
specification.
-* Marking global variables with `declare target`'s `link` clause and mapping
-the data to the device data environment utilising `declare target`. This may
-not work for all types yet, but for scalars and arrays of scalars, it
+* Marking global variables with `declare target`'s `link` clause and mapping
+the data to the device data environment utilising `declare target`. This may
+not work for all types yet, but for scalars and arrays of scalars, it
should.
Doesn't work for, or needs further testing for:
-* Marking the following types with `declare target link` (needs further
+* Marking the following types with `declare target link` (needs further
testing):
* Descriptor based types, e.g. pointers/allocatables.
* Derived types.
* Members of derived types (use-case needs legality checking with OpenMP
specification).
* Marking global variables with `declare target`'s `to` clause. A lot of the
-lowering should exist, but it needs further testing and likely some further
+lowering should exist, but it needs further testing and likely some further
changes to fully function.
diff --git a/flang/docs/OpenMP-descriptor-management.md b/flang/docs/OpenMP-descriptor-management.md
index d0764c2df4642..d2a3c93fa1647 100644
--- a/flang/docs/OpenMP-descriptor-management.md
+++ b/flang/docs/OpenMP-descriptor-management.md
@@ -15,11 +15,11 @@ local:
## Details
-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
+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
+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++
@@ -34,7 +34,7 @@ int main() {
mock_descriptor data;
#pragma omp target map(tofrom: data, data.pointer[:upper_bound])
{
- do something...
+ do something...
}
return 0;
@@ -42,19 +42,19 @@ mock_descriptor data;
```
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
+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,
-`MapInfoFinalizationPass` (Optimizer/OpenMP/MapInfoFinalization.cpp) will expand the
-`omp.MapInfoOp`'s containing descriptors (which currently will be a `BoxType` or `BoxAddrOp`) into multiple
+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,
+`MapInfoFinalizationPass` (Optimizer/OpenMP/MapInfoFinalization.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
+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.TargetDataOp` etc.) map operand list and in cases where the owning
operation is `IsolatedFromAbove`, it also inserts them as `BlockArgs` to canonicalize the mappings and
@@ -85,26 +85,26 @@ omp.target map_entries(%13 -> %arg1, %14 -> %arg2, %15 -> %arg3 : !fir.llvm_ptr<
```
In later stages of the compilation flow when the OpenMP dialect is being lowered to LLVM-IR these descriptor
-mappings are treated as if they were structure mappings with explicit member maps on the same directive as
-their parent was mapped.
+mappings are treated as if they were structure mappings with explicit member maps on the same directive as
+their parent was mapped.
-This implementation utilises the member field of the `map_info` operation to indicate that the pointer
-descriptor elements which are contained in their own `map_info` operation are part of their respective
+This implementation utilises the member field of the `map_info` operation to indicate that the pointer
+descriptor elements which are contained in their own `map_info` operation are part of their respective
parent descriptor. This allows the descriptor containing the descriptor pointer member to be mapped
as a composite entity during lowering, with the correct mappings being generated to tie them together,
allowing the OpenMP runtime to map them correctly, attaching the pointer member to the parent
-structure so it can be accessed during execution. If we opt to not treat the descriptor as a single
+structure so it can be accessed during execution. If we opt to not treat the descriptor as a single
entity we have issues with the member being correctly attached to the parent and being accessible,
this can cause runtime segfaults on the device when we try to access the data through the parent. It
-may be possible to avoid this member mapping, treating them as individual entities, but treating a
-composite mapping as an individual mapping could lead to problems such as the runtime taking
-liberties with the mapping it usually wouldn't if it knew they were linked, we would also have to
+may be possible to avoid this member mapping, treating them as individual entities, but treating a
+composite mapping as an individual mapping could lead to problems such as the runtime taking
+liberties with the mapping it usually wouldn't if it knew they were linked, we would also have to
be careful to maintian the correct order of mappings as we lower, if we misorder the maps, it'd be
possible to overwrite already written data, e.g. if we write the descriptor data pointer first, and
-then the containing descriptor, we would overwrite the descriptor data pointer with the incorrect
+then the containing descriptor, we would overwrite the descriptor data pointer with the incorrect
address.
-This method is generic in the sense that the OpenMP dialect doesn't need to understand that it is mapping a
+This method is generic in the sense that the OpenMP dialect doesn't need to understand that it is mapping a
Fortran type containing a descriptor, it just thinks it's a record type from either Fortran or C++. However,
it is a little rigid in how the descriptor mappings are handled as there is no specialisation or possibility
to specialise the mappings for possible edge cases without polluting the dialect or lowering with further
@@ -112,22 +112,22 @@ knowledge of Fortran and the FIR dialect.
## Differences from OpenACC
-The descriptor mapping for OpenMP currently works differently to the planned direction for OpenACC, however,
-it is possible and would likely be ideal to align the method with OpenACC in the future.
+The descriptor mapping for OpenMP currently works differently to the planned direction for OpenACC, however,
+it is possible and would likely be ideal to align the method with OpenACC in the future.
Currently the OpenMP specification is less descriptive and has less stringent rules around descriptor based
-types so does not require as complex a set of descriptor management rules as OpenACC (although, in certain
-cases for the interim adopting OpenACC's rules where it makes sense could be useful). To handle the more
-complex descriptor mapping rules OpenACC has opted to utilise a more runtime oriented approach, where
-specialized runtime functions for handling descriptor mapping for OpenACC are created and these runtime
-function handles are attatched to a special OpenACC dialect operation. When this operation is lowered it
-will lower to the attatched OpenACC descriptor mapping runtime function. This sounds like it will work
-(no implementation yet) similarly to some of the existing HLFIR operations which optionally lower to
-Fortran runtime calls.
+types so does not require as complex a set of descriptor management rules as OpenACC (although, in certain
+cases for the interim adopting OpenACC's rules where it makes sense could be useful). To handle the more
+complex descriptor mapping rules OpenACC has opted to utilise a more runtime oriented approach, where
+specialized runtime functions for handling descriptor mapping for OpenACC are created and these runtime
+function handles are attatched to a special OpenACC dialect operation. When this operation is lowered it
+will lower to the attatched OpenACC descriptor mapping runtime function. This sounds like it will work
+(no implementation yet) similarly to some of the existing HLFIR operations which optionally lower to
+Fortran runtime calls.
This methodology described by OpenACC which utilises runtime functions to handle specialised mappings allows
more flexibility as a significant amount of the mapping logic can be moved into the runtime from the compiler.
It also allows specialisation of the mapping for fortran specific types. This may be a desireable approach
-to take for OpenMP in the future, in particular if we find need to specialise mapping further for
+to take for OpenMP in the future, in particular if we find need to specialise mapping further for
descriptors or other Fortran types. However, for the moment the currently chosen implementation for OpenMP
appears sufficient as far as the OpenMP specification and current testing can show.
diff --git a/flang/docs/OpenMP-semantics.md b/flang/docs/OpenMP-semantics.md
index 9ba7595ce6130..de85f98292d10 100644
--- a/flang/docs/OpenMP-semantics.md
+++ b/flang/docs/OpenMP-semantics.md
@@ -1,9 +1,9 @@
-<!--===- docs/OpenMP-semantics.md
-
+<!--===- docs/OpenMP-semantics.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 Semantic Analysis
diff --git a/flang/docs/OptionComparison.md b/flang/docs/OptionComparison.md
index 776a6df3f4c4d..05000e5a7596f 100644
--- a/flang/docs/OptionComparison.md
+++ b/flang/docs/OptionComparison.md
@@ -1,16 +1,16 @@
-<!--===- docs/OptionComparison.md
-
+<!--===- docs/OptionComparison.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
-
+
-->
# Compiler options comparison
This document catalogs the options processed by Flang's peers/competitors. Much of the document is taken up by a set of tables that list the options categorized into different topics. Some of the table headings link to more information about the contents of the tables. For example, the table on **Standards conformance** options links to <a href="#standards">notes on Standards conformance</a>.
-**There's also important information in the ___[Appendix section](#appendix)___ near the end of the document on how this data was gathered and what ___is___ and ___is not___ included in this document.**
+**There's also important information in the ___[Appendix section](#appendix)___ near the end of the document on how this data was gathered and what ___is___ and ___is not___ included in this document.**
Note that compilers may support language features without having an option for them. Such cases are frequently, but not always noted in this document.
@@ -281,7 +281,7 @@ Tf filename
</td>
<td>eQ
</td>
- <td>N/A
+ <td>N/A
</td>
<td>N/A
</td>
@@ -428,7 +428,7 @@ qonetrip
</td>
<td>Monetrip
</td>
- <td>N/A
+ <td>N/A
</td>
</tr>
<tr>
@@ -1187,16 +1187,16 @@ Mcuda
## Notes
-**<a name="standards"></a>Standards conformance:**
+**<a name="standards"></a>Standards conformance:**
All conformance options are similar -- they issue warnings if non-standard features are used. All defaults are to allow extensions without warnings. The GNU, IBM, and Intel compilers allow multiple standard levels to be specified.
-* **Cray**: The capital "-eN" option specifies to issue error messages for non-compliance rather than warnings.
+* **Cray**: The capital "-eN" option specifies to issue error messages for non-compliance rather than warnings.
* **GNU:** The "std=_level_" option specifies the standard to which the program is expected to conform. The default value for std is 'gnu', which specifies a superset of the latest Fortran standard that includes all of the extensions supported by GNU Fortran, although warnings will be given for obsolete extensions not recommended for use in new code. The 'legacy' value is equivalent but without the warnings for obsolete extensions. The 'f95', 'f2003', 'f2008', and 'f2018' values specify strict conformance to the respective standards. Errors are given for all extensions beyond the relevant language standard, and warnings are given for the Fortran 77 features that are permitted but obsolescent in later standards. '-std=f2008ts' allows the Fortran 2008 standard including the additions of the Technical Specification (TS) 29113 on Further Interoperability of Fortran with C and TS 18508 on Additional Parallel Features in Fortran. Values for "_level_" are f_95, f2003, f2008, f2008ts, f2018, gnu,_ and _legacy._
-**<a name="source"></a>Source format:**
+**<a name="source"></a>Source format:**
**Fixed or free source:** Cray, IBM, and Intel default the source format based on the source file suffix as follows:
@@ -1220,7 +1220,7 @@ IBM Fortran's options allow the source line length to be specified with the opti
* **GNU:** For both "ffixed-line-length-_n_" and "ffree-line-length-_n_" options, characters are ignored after the specified length. The default for fixed is 72. The default for free is 132. For free, you can specify 'none' as the length, which means that all characters in the line are meaningful.
* **IBM:** For **fixed**, the default is 72. For **free**, there's no default, but the maximum length for either form is 132.
* **Intel:** The default is 72 for **fixed** and 132 for **free**.
-* **PGI, Classic Flang:**
+* **PGI, Classic Flang:**
* in free form, it is an error if the line is longer than 1000 characters
* in fixed form by default, characters after column 72 are ignored
* in fixed form with -Mextend, characters after column 132 are ignored
diff --git a/flang/docs/Overview.md b/flang/docs/Overview.md
index dfb4d89264a75..77ed21c021f8a 100644
--- a/flang/docs/Overview.md
+++ b/flang/docs/Overview.md
@@ -1,9 +1,9 @@
-<!--===- docs/Overview.md
-
+<!--===- docs/Overview.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
-
+
-->
# Overview of Compiler Phases
@@ -13,7 +13,7 @@
local:
---
```
-The Flang compiler transforms Fortran source code into an executable file.
+The Flang compiler transforms Fortran source code into an executable file.
This transformation proceeds in three high level phases -- analysis, lowering,
and code generation/linking.
@@ -64,7 +64,7 @@ See [Preprocessing.md](Preprocessing.md).
**Entry point:** `parser::Parsing::Prescan`
-**Commands:**
+**Commands:**
- `flang -fc1 -E src.f90` dumps the cooked character stream
- `flang -fc1 -fdebug-dump-provenance src.f90` dumps provenance
information
@@ -90,7 +90,7 @@ representing a syntactically correct program, rooted at the program unit. See:
**Input:** the parse tree, the cooked character stream, and provenance
information
-**Output:**
+**Output:**
* a symbol table
* modified parse tree
* module files, (see: [ModFiles.md](ModFiles.md))
@@ -101,15 +101,15 @@ information
**Entry point:** `semantics::Semantics::Perform`
For more detail on semantic analysis, see: [Semantics.md](Semantics.md).
-Semantic processing performs several tasks:
+Semantic processing performs several tasks:
* validates labels, see: [LabelResolution.md](LabelResolution.md).
-* canonicalizes DO statements,
+* canonicalizes DO statements,
* canonicalizes OpenACC and OpenMP code
* resolves names, building a tree of scopes and symbols
* rewrites the parse tree to correct parsing mistakes (when needed) once semantic information is available to clarify the program's meaning
* checks the validity of declarations
* analyzes expressions and statements, emitting error messages where appropriate
-* creates module files if the source code contains modules,
+* creates module files if the source code contains modules,
see [ModFiles.md](ModFiles.md).
In the course of semantic analysis, the compiler:
@@ -132,7 +132,7 @@ produces LLVM IR.
### Create the lowering bridge
-**Inputs:**
+**Inputs:**
- the parse tree
- the symbol table
- The default KINDs for intrinsic types (specified by default or command line option)
diff --git a/flang/docs/ParallelMultiImageFortranRuntime.md b/flang/docs/ParallelMultiImageFortranRuntime.md
index 002959a39e7cf..6dd21a420d787 100644
--- a/flang/docs/ParallelMultiImageFortranRuntime.md
+++ b/flang/docs/ParallelMultiImageFortranRuntime.md
@@ -1,18 +1,17 @@
<!--===- docs/ParallelMultiImageFortranRuntime.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
-
+
-->
# Multi-Image Parallel Fortran Runtime
-The Parallel Runtime Interface for Fortran (PRIF) defines an
-interface designed for LLVM Flang to target implementations of
+The Parallel Runtime Interface for Fortran (PRIF) defines an
+interface designed for LLVM Flang to target implementations of
Fortran's multi-image parallel features.
The current revision of the PRIF specification is here:
<https://doi.org/10.25344/S46S3W>
-
diff --git a/flang/docs/ParameterizedDerivedTypes.md b/flang/docs/ParameterizedDerivedTypes.md
index 0ed9f8b494e78..0e6b62fffe7e9 100644
--- a/flang/docs/ParameterizedDerivedTypes.md
+++ b/flang/docs/ParameterizedDerivedTypes.md
@@ -786,7 +786,7 @@ subroutine foo(x, n)
! type parameters of `x(n)%pdt_component` are not propagated from the caller.
! A descriptor local to this function is created to pass the array section
- ! in bar.
+ ! in bar.
call bar(x%pdt_component)
end subroutine
diff --git a/flang/docs/ParserCombinators.md b/flang/docs/ParserCombinators.md
index 72cc4ba00307b..c594d6d96cce3 100644
--- a/flang/docs/ParserCombinators.md
+++ b/flang/docs/ParserCombinators.md
@@ -1,9 +1,9 @@
-<!--===- docs/ParserCombinators.md
-
+<!--===- docs/ParserCombinators.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
-
+
-->
# Parser Combinators
diff --git a/flang/docs/Parsing.md b/flang/docs/Parsing.md
index ccc73a1d98524..eae92519c6004 100644
--- a/flang/docs/Parsing.md
+++ b/flang/docs/Parsing.md
@@ -1,9 +1,9 @@
-<!--===- docs/Parsing.md
-
+<!--===- docs/Parsing.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
-
+
-->
# The Flang Parser
@@ -182,7 +182,7 @@ Parse tree entities should be viewed as values, not objects; their
addresses should not be abused for purposes of identification. They are
assembled with C++ move semantics during parse tree construction.
Their default and copy constructors are deliberately deleted in their
-declarations.
+declarations.
The std::list<> data type is used in the parse tree to reliably store pointers
to other relevant entries in the tree. Since the tree lists are moved and
diff --git a/flang/docs/Preprocessing.md b/flang/docs/Preprocessing.md
index 5df07bc125686..e167922319785 100644
--- a/flang/docs/Preprocessing.md
+++ b/flang/docs/Preprocessing.md
@@ -1,9 +1,9 @@
-<!--===- docs/Preprocessing.md
-
+<!--===- docs/Preprocessing.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
-
+
-->
# Fortran Preprocessing
diff --git a/flang/docs/PullRequestChecklist.md b/flang/docs/PullRequestChecklist.md
index f1cf498a1624c..5887def177fb6 100644
--- a/flang/docs/PullRequestChecklist.md
+++ b/flang/docs/PullRequestChecklist.md
@@ -1,4 +1,4 @@
-<!--===- docs/PullRequestChecklist.md
+<!--===- docs/PullRequestChecklist.md
Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
See https://llvm.org/LICENSE.txt for license information.
diff --git a/flang/docs/RuntimeDescriptor.md b/flang/docs/RuntimeDescriptor.md
index c90bad5d18ee1..a2a701a8c02bd 100644
--- a/flang/docs/RuntimeDescriptor.md
+++ b/flang/docs/RuntimeDescriptor.md
@@ -1,9 +1,9 @@
-<!--===- docs/RuntimeDescriptor.md
-
+<!--===- docs/RuntimeDescriptor.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
-
+
-->
# Runtime Descriptors
diff --git a/flang/docs/RuntimeEnvironment.md b/flang/docs/RuntimeEnvironment.md
index ccc401c51e8a3..e762ce770fd83 100644
--- a/flang/docs/RuntimeEnvironment.md
+++ b/flang/docs/RuntimeEnvironment.md
@@ -1,9 +1,9 @@
-<!--===- docs/RuntimeEnvironment.md
-
+<!--===- docs/RuntimeEnvironment.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
-
+
-->
```{contents}
diff --git a/flang/docs/RuntimeTypeInfo.md b/flang/docs/RuntimeTypeInfo.md
index 16fff4dd563ee..c41a9975c3b9f 100644
--- a/flang/docs/RuntimeTypeInfo.md
+++ b/flang/docs/RuntimeTypeInfo.md
@@ -1,9 +1,9 @@
-<!--===- docs/RuntimeTypeInfo.md
-
+<!--===- docs/RuntimeTypeInfo.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
-
+
-->
# The derived type runtime information table
diff --git a/flang/docs/Semantics.md b/flang/docs/Semantics.md
index 0fc1ebe4cff1d..6427e2c778185 100644
--- a/flang/docs/Semantics.md
+++ b/flang/docs/Semantics.md
@@ -1,9 +1,9 @@
-<!--===- docs/Semantics.md
-
+<!--===- docs/Semantics.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
-
+
-->
# Semantic Analysis
diff --git a/flang/docs/Unsigned.md b/flang/docs/Unsigned.md
index 5815822cd53ad..be99ba9b616ef 100644
--- a/flang/docs/Unsigned.md
+++ b/flang/docs/Unsigned.md
@@ -1,9 +1,9 @@
-<!--===- docs/Unsigned.md
-
+<!--===- docs/Unsigned.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
-
+
-->
# Flang support for UNSIGNED type
More information about the flang-commits
mailing list