Skip to content

Commit be14124

Browse files
agozillonagozillon
committed
[Flang][OpenMP] Initial mapping of Fortran pointers, allocatables and other descriptor types for target devices
This patch seeks to add an initial lowering for pointers and allocatable variables captured by implicit and explicit map in Flang OpenMP for Target operations that take map clauses e.g. Target, Target Update. Target Exit/Enter etc. Currently this is done by treating the type that lowers to a descriptor (allocatable/pointer/assumed shape) as a map of a record type (e.g. a structure) as that's effectively what descriptor types lower to in LLVM-IR and what they're represented as in the Fortran runtime (written in C/C++). The descriptor effectively lowers to a strucutre containing scalar and array elements that represent various aspects of the underlying data being mapped (lower bound, upper bound, extent being the main ones of interest in most cases) and a pointer to the allocated data. In this current iteration of the mapping we map the structure in it's entirety and then attach the underlying data pointer and map the data to the device, this allows most of the required data to be resident on the device for use. Currently we do not support the addendum (another block of pointer data), but it shouldn't be too difficult to extend this to support it. The MapInfoOp generation for descriptor types is primarily handled in an optimisation pass, where it expands BoxType (descriptor types) map captures into two maps, one for the structure (scalar elements) and the other for the pointer data (base address) and links them in a Parent <-> Child relationship. The later lowering processes will then treat them as a conjoined structure with a pointer member map.
1 parent bc82d1a commit be14124

36 files changed

+1657
-95
lines changed

‎flang/docs/OpenMP-descriptor-management.md‎

Lines changed: 125 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
1+
<!--===- docs/OpenMP-descriptor-management.md
2+
3+
Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
See https://llvm.org/LICENSE.txt for license information.
5+
SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
7+
-->
8+
9+
# OpenMP dialect: Fortran descriptor type mapping for offload
10+
11+
The initial method for mapping Fortran types tied to descriptors for OpenMP offloading is to treat these types
12+
as a special case of OpenMP record type (C/C++ structure/class, Fortran derived type etc.) mapping as far as the
13+
runtime is concerned. Where the box (descriptor information) is the holding container and the underlying
14+
data pointer is contained within the container, and we must generate explicit maps for both the pointer member and
15+
the container. As an example, a small C++ program that is equivalent to the concept described, with the
16+
`mock_descriptor` class being representative of the class utilised for descriptors in Clang:
17+
18+
```C++
19+
struct mock_descriptor {
20+
long int x;
21+
std::byte x1, x2, x3, x4;
22+
void *pointer;
23+
long int lx[1][3];
24+
};
25+
26+
int main() {
27+
mock_descriptor data;
28+
#pragma omp target map(tofrom: data, data.pointer[:upper_bound])
29+
{
30+
do something...
31+
}
32+
33+
return 0;
34+
}
35+
```
36+
37+
In the above, we have to map both the containing structure, with its non-pointer members and the
38+
data pointed to by the pointer contained within the structure to appropriately access the data. This
39+
is effectively what is done with descriptor types for the time being. Other pointers that are part
40+
of the descriptor container such as the addendum should also be treated as the data pointer is
41+
treated.
42+
43+
Currently, Flang will lower these descriptor types in the OpenMP lowering (lower/OpenMP.cpp) similarly
44+
to all other map types, generating an omp.MapInfoOp containing relevant information required for lowering
45+
the OpenMP dialect to LLVM-IR during the final stages of the MLIR lowering. However, after
46+
the lowering to FIR/HLFIR has been performed an OpenMP dialect specific pass for Fortran,
47+
`OMPDescriptorMapInfoGenPass` (Optimizer/OMPDescriptorMapInfoGen.cpp) will expand the
48+
`omp.MapInfoOp`'s containing descriptors (which currently will be a `BoxType` or `BoxAddrOp`) into multiple
49+
mappings, with one extra per pointer member in the descriptor that is supported on top of the original
50+
descriptor map operation. These pointers members are linked to the parent descriptor by adding them to
51+
the member field of the original descriptor map operation, they are then inserted into the relevant map
52+
owning operation's (`omp.TargetOp`, `omp.DataOp` etc.) map operand list and in cases where the owning operation
53+
is `IsolatedFromAbove`, it also inserts them as `BlockArgs` to canonicalize the mappings and simplify lowering.
54+
55+
An example transformation by the `OMPDescriptorMapInfoGenPass`:
56+
57+
```
58+
59+
...
60+
%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"}
61+
...
62+
omp.target map_entries(%12 -> %arg1, %13 -> %arg2 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<i32>) {
63+
^bb0(%arg1: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, %arg2: !fir.ref<i32>):
64+
...
65+
66+
====>
67+
68+
...
69+
%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>>>
70+
%13 = omp.map_info var_ptr(%1#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.array<?xi32>) var_ptr_ptr(%12 : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%11) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
71+
%14 = 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) members(%13 : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> {name = "arg_alloc"}
72+
...
73+
omp.target map_entries(%13 -> %arg1, %14 -> %arg2, %15 -> %arg3 : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<i32>) {
74+
^bb0(%arg1: !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, %arg2: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, %arg3: !fir.ref<i32>):
75+
...
76+
77+
```
78+
79+
In later stages of the compilation flow when the OpenMP dialect is being lowered to LLVM-IR these descriptor
80+
mappings are treated as if they were structure mappings with explicit member maps on the same directive as
81+
their parent was mapped.
82+
83+
This implementation utilises the member field of the `map_info` operation to indicate that the pointer
84+
descriptor elements which are contained in their own `map_info` operation are part of their respective
85+
parent descriptor. This allows the descriptor containing the descriptor pointer member to be mapped
86+
as a composite entity during lowering, with the correct mappings being generated to tie them together,
87+
allowing the OpenMP runtime to map them correctly, attaching the pointer member to the parent
88+
structure so it can be accessed during execution. If we opt to not treat the descriptor as a single
89+
entity we have issues with the member being correctly attached to the parent and being accessible,
90+
this can cause runtime segfaults on the device when we try to access the data through the parent. It
91+
may be possible to avoid this member mapping, treating them as individual entities, but treating a
92+
composite mapping as an individual mapping could lead to problems such as the runtime taking
93+
liberties with the mapping it usually wouldn't if it knew they were linked, we would also have to
94+
be careful to maintian the correct order of mappings as we lower, if we misorder the maps, it'd be
95+
possible to overwrite already written data, e.g. if we write the descriptor data pointer first, and
96+
then the containing descriptor, we would overwrite the descriptor data pointer with the incorrect
97+
address.
98+
99+
This method is generic in the sense that the OpenMP dialect doesn't need to understand that it is mapping a
100+
Fortran type containing a descriptor, it just thinks it's a record type from either Fortran or C++. However,
101+
it is a little rigid in how the descriptor mappings are handled as there is no specialisation or possibility
102+
to specialise the mappings for possible edge cases without polluting the dialect or lowering with further
103+
knowledge of Fortran and the FIR dialect.
104+
105+
# OpenMP dialect differences from OpenACC dialect
106+
107+
The descriptor mapping for OpenMP currently works differently to the planned direction for OpenACC, however,
108+
it is possible and would likely be ideal to align the method with OpenACC in the future.
109+
110+
Currently the OpenMP specification is less descriptive and has less stringent rules around descriptor based
111+
types so does not require as complex a set of descriptor management rules as OpenACC (although, in certain
112+
cases for the interim adopting OpenACC's rules where it makes sense could be useful). To handle the more
113+
complex descriptor mapping rules OpenACC has opted to utilise a more runtime oriented approach, where
114+
specialized runtime functions for handling descriptor mapping for OpenACC are created and these runtime
115+
function handles are attatched to a special OpenACC dialect operation. When this operation is lowered it
116+
will lower to the attatched OpenACC descriptor mapping runtime function. This sounds like it will work
117+
(no implementation yet) similarly to some of the existing HLFIR operations which optionally lower to
118+
Fortran runtime calls.
119+
120+
This methodology described by OpenACC which utilises runtime functions to handle specialised mappings allows
121+
more flexibility as a significant amount of the mapping logic can be moved into the runtime from the compiler.
122+
It also allows specialisation of the mapping for fortran specific types. This may be a desireable approach
123+
to take for OpenMP in the future, in particular if we find need to specialise mapping further for
124+
descriptors or other Fortran types. However, for the moment the currently chosen implementation for OpenMP
125+
appears sufficient as far as the OpenMP specification and current testing can show.

‎flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
//===------- Optimizer/CodeGen/CodeGenOpenMP.h - OpenMP codegen -*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#ifndef FORTRAN_OPTIMIZER_CODEGEN_CODEGENOPENMP_H
10+
#define FORTRAN_OPTIMIZER_CODEGEN_CODEGENOPENMP_H
11+
12+
#include "mlir/Pass/Pass.h"
13+
#include "mlir/Pass/PassRegistry.h"
14+
15+
namespace fir {
16+
class LLVMTypeConverter;
17+
18+
/// Specialised conversion patterns of OpenMP operations for FIR to LLVM
19+
/// dialect, utilised in cases where the default OpenMP dialect handling cannot
20+
/// handle all cases for intermingled fir types and operations.
21+
void populateOpenMPFIRToLLVMConversionPatterns(
22+
LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns);
23+
24+
} // namespace fir
25+
26+
#endif // FORTRAN_OPTIMIZER_CODEGEN_CODEGENOPENMP_H

‎flang/include/flang/Optimizer/Dialect/FIRType.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -321,6 +321,9 @@ bool isBoxNone(mlir::Type ty);
321321
/// e.g. !fir.box<!fir.type<derived>>
322322
bool isBoxedRecordType(mlir::Type ty);
323323

324+
/// Return true iff `ty` is a type that contains descriptor information.
325+
bool isTypeWithDescriptor(mlir::Type ty);
326+
324327
/// Return true iff `ty` is a scalar boxed record type.
325328
/// e.g. !fir.box<!fir.type<derived>>
326329
/// !fir.box<!fir.heap<!fir.type<derived>>>

‎flang/include/flang/Optimizer/Transforms/Passes.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -76,6 +76,7 @@ std::unique_ptr<mlir::Pass>
7676
createAlgebraicSimplificationPass(const mlir::GreedyRewriteConfig &config);
7777
std::unique_ptr<mlir::Pass> createPolymorphicOpConversionPass();
7878

79+
std::unique_ptr<mlir::Pass> createOMPDescriptorMapInfoGenPass();
7980
std::unique_ptr<mlir::Pass> createOMPFunctionFilteringPass();
8081
std::unique_ptr<mlir::OperationPass<mlir::ModuleOp>>
8182
createOMPMarkDeclareTargetPass();

‎flang/include/flang/Optimizer/Transforms/Passes.td‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -318,6 +318,18 @@ def LoopVersioning : Pass<"loop-versioning", "mlir::func::FuncOp"> {
318318
let dependentDialects = [ "fir::FIROpsDialect" ];
319319
}
320320

321+
def OMPDescriptorMapInfoGenPass
322+
: Pass<"omp-descriptor-map-info-gen", "mlir::func::FuncOp"> {
323+
let summary = "expands OpenMP MapInfo operations containing descriptors";
324+
let description = [{
325+
Expands MapInfo operations containing descriptor types into multiple
326+
MapInfo's for each pointer element in the descriptor that requires
327+
explicit individual mapping by the OpenMP runtime.
328+
}];
329+
let constructor = "::fir::createOMPDescriptorMapInfoGenPass()";
330+
let dependentDialects = ["mlir::omp::OpenMPDialect"];
331+
}
332+
321333
def OMPMarkDeclareTargetPass
322334
: Pass<"omp-mark-declare-target", "mlir::ModuleOp"> {
323335
let summary = "Marks all functions called by an OpenMP declare target function as declare target";

‎flang/include/flang/Tools/CLOptions.inc‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -274,6 +274,7 @@ inline void createHLFIRToFIRPassPipeline(
274274
/// rather than the host device.
275275
inline void createOpenMPFIRPassPipeline(
276276
mlir::PassManager &pm, bool isTargetDevice) {
277+
pm.addPass(fir::createOMPDescriptorMapInfoGenPass());
277278
pm.addPass(fir::createOMPMarkDeclareTargetPass());
278279
if (isTargetDevice)
279280
pm.addPass(fir::createOMPFunctionFilteringPass());

‎flang/lib/Lower/OpenMP.cpp‎

Lines changed: 48 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -1821,27 +1821,25 @@ bool ClauseProcessor::processLink(
18211821

18221822
static mlir::omp::MapInfoOp
18231823
createMapInfoOp(fir::FirOpBuilder &builder, mlir::Location loc,
1824-
mlir::Value baseAddr, std::stringstream &name,
1825-
mlir::SmallVector<mlir::Value> bounds, uint64_t mapType,
1826-
mlir::omp::VariableCaptureKind mapCaptureType,
1827-
mlir::Type retTy) {
1828-
mlir::Value varPtr, varPtrPtr;
1829-
mlir::TypeAttr varType;
1830-
1824+
mlir::Value baseAddr, mlir::Value varPtrPtr, std::string name,
1825+
mlir::SmallVector<mlir::Value> bounds,
1826+
mlir::SmallVector<mlir::Value> members, uint64_t mapType,
1827+
mlir::omp::VariableCaptureKind mapCaptureType, mlir::Type retTy,
1828+
bool isVal = false) {
18311829
if (auto boxTy = baseAddr.getType().dyn_cast<fir::BaseBoxType>()) {
18321830
baseAddr = builder.create<fir::BoxAddrOp>(loc, baseAddr);
18331831
retTy = baseAddr.getType();
18341832
}
18351833

1836-
varPtr = baseAddr;
1837-
varType = mlir::TypeAttr::get(
1834+
mlir::TypeAttr varType = mlir::TypeAttr::get(
18381835
llvm::cast<mlir::omp::PointerLikeType>(retTy).getElementType());
18391836

18401837
mlir::omp::MapInfoOp op = builder.create<mlir::omp::MapInfoOp>(
1841-
loc, retTy, varPtr, varType, varPtrPtr, bounds,
1838+
loc, retTy, baseAddr, varType, varPtrPtr, members, bounds,
18421839
builder.getIntegerAttr(builder.getIntegerType(64, false), mapType),
18431840
builder.getAttr<mlir::omp::VariableCaptureKindAttr>(mapCaptureType),
1844-
builder.getStringAttr(name.str()));
1841+
builder.getStringAttr(name));
1842+
18451843
return op;
18461844
}
18471845

@@ -1904,28 +1902,37 @@ bool ClauseProcessor::processMap(
19041902
std::get<Fortran::parser::OmpObjectList>(mapClause->v.t).v) {
19051903
llvm::SmallVector<mlir::Value> bounds;
19061904
std::stringstream asFortran;
1905+
19071906
Fortran::lower::AddrAndBoundsInfo info =
19081907
Fortran::lower::gatherDataOperandAddrAndBounds<
19091908
Fortran::parser::OmpObject, mlir::omp::DataBoundsOp,
19101909
mlir::omp::DataBoundsType>(
19111910
converter, firOpBuilder, semanticsContext, stmtCtx, ompObject,
19121911
clauseLocation, asFortran, bounds, treatIndexAsSection);
19131912

1913+
auto origSymbol =
1914+
converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
1915+
mlir::Value symAddr = info.addr;
1916+
if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType()))
1917+
symAddr = origSymbol;
1918+
19141919
// Explicit map captures are captured ByRef by default,
19151920
// optimisation passes may alter this to ByCopy or other capture
19161921
// types to optimise
19171922
mlir::Value mapOp = createMapInfoOp(
1918-
firOpBuilder, clauseLocation, info.addr, asFortran, bounds,
1923+
firOpBuilder, clauseLocation, symAddr, mlir::Value{},
1924+
asFortran.str(), bounds, {},
19191925
static_cast<
19201926
std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
19211927
mapTypeBits),
1922-
mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
1928+
mlir::omp::VariableCaptureKind::ByRef, symAddr.getType());
19231929

19241930
mapOperands.push_back(mapOp);
19251931
if (mapSymTypes)
1926-
mapSymTypes->push_back(info.addr.getType());
1932+
mapSymTypes->push_back(symAddr.getType());
19271933
if (mapSymLocs)
1928-
mapSymLocs->push_back(info.addr.getLoc());
1934+
mapSymLocs->push_back(symAddr.getLoc());
1935+
19291936
if (mapSymbols)
19301937
mapSymbols->push_back(getOmpObjectSymbol(ompObject));
19311938
}
@@ -2032,12 +2039,22 @@ bool ClauseProcessor::processMotionClauses(
20322039
converter, firOpBuilder, semanticsContext, stmtCtx, ompObject,
20332040
clauseLocation, asFortran, bounds, treatIndexAsSection);
20342041

2042+
auto origSymbol =
2043+
converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
2044+
mlir::Value symAddr = info.addr;
2045+
if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType()))
2046+
symAddr = origSymbol;
2047+
2048+
// Explicit map captures are captured ByRef by default,
2049+
// optimisation passes may alter this to ByCopy or other capture
2050+
// types to optimise
20352051
mlir::Value mapOp = createMapInfoOp(
2036-
firOpBuilder, clauseLocation, info.addr, asFortran, bounds,
2052+
firOpBuilder, clauseLocation, symAddr, mlir::Value{},
2053+
asFortran.str(), bounds, {},
20372054
static_cast<
20382055
std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
20392056
mapTypeBits),
2040-
mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
2057+
mlir::omp::VariableCaptureKind::ByRef, symAddr.getType());
20412058

20422059
mapOperands.push_back(mapOp);
20432060
}
@@ -2812,7 +2829,8 @@ static void genBodyOfTargetOp(
28122829
std::stringstream name;
28132830
firOpBuilder.setInsertionPoint(targetOp);
28142831
mlir::Value mapOp = createMapInfoOp(
2815-
firOpBuilder, copyVal.getLoc(), copyVal, name, bounds,
2832+
firOpBuilder, copyVal.getLoc(), copyVal, mlir::Value{}, name.str(),
2833+
bounds, llvm::SmallVector<mlir::Value>{},
28162834
static_cast<
28172835
std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
28182836
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT),
@@ -2934,18 +2952,21 @@ genTargetOp(Fortran::lower::AbstractConverter &converter,
29342952
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT;
29352953
mlir::omp::VariableCaptureKind captureKind =
29362954
mlir::omp::VariableCaptureKind::ByRef;
2937-
if (auto refType = baseOp.getType().dyn_cast<fir::ReferenceType>()) {
2938-
auto eleType = refType.getElementType();
2939-
if (fir::isa_trivial(eleType) || fir::isa_char(eleType)) {
2940-
captureKind = mlir::omp::VariableCaptureKind::ByCopy;
2941-
} else if (!fir::isa_builtin_cptr_type(eleType)) {
2942-
mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TO;
2943-
mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_FROM;
2944-
}
2955+
2956+
mlir::Type eleType = baseOp.getType();
2957+
if (auto refType = baseOp.getType().dyn_cast<fir::ReferenceType>())
2958+
eleType = refType.getElementType();
2959+
2960+
if (fir::isa_trivial(eleType) || fir::isa_char(eleType)) {
2961+
captureKind = mlir::omp::VariableCaptureKind::ByCopy;
2962+
} else if (!fir::isa_builtin_cptr_type(eleType)) {
2963+
mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TO;
2964+
mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_FROM;
29452965
}
29462966

29472967
mlir::Value mapOp = createMapInfoOp(
2948-
converter.getFirOpBuilder(), baseOp.getLoc(), baseOp, name, bounds,
2968+
converter.getFirOpBuilder(), baseOp.getLoc(), baseOp, mlir::Value{},
2969+
name.str(), bounds, {},
29492970
static_cast<
29502971
std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
29512972
mapFlag),

‎flang/lib/Optimizer/CodeGen/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@ add_flang_library(FIRCodeGen
22
BoxedProcedure.cpp
33
CGOps.cpp
44
CodeGen.cpp
5+
CodeGenOpenMP.cpp
56
PreCGRewrite.cpp
67
TBAABuilder.cpp
78
Target.cpp

‎flang/lib/Optimizer/CodeGen/CodeGen.cpp‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@
1313
#include "flang/Optimizer/CodeGen/CodeGen.h"
1414

1515
#include "CGOps.h"
16+
#include "flang/Optimizer/CodeGen/CodeGenOpenMP.h"
1617
#include "flang/Optimizer/Dialect/FIRAttr.h"
1718
#include "flang/Optimizer/Dialect/FIROps.h"
1819
#include "flang/Optimizer/Dialect/FIRType.h"
@@ -3959,6 +3960,11 @@ class FIRToLLVMLowering
39593960
mlir::populateMathToLibmConversionPatterns(pattern);
39603961
mlir::populateComplexToLLVMConversionPatterns(typeConverter, pattern);
39613962
mlir::populateVectorToLLVMConversionPatterns(typeConverter, pattern);
3963+
3964+
// Flang specific overloads for OpenMP operations, to allow for special
3965+
// handling of things like Box types.
3966+
fir::populateOpenMPFIRToLLVMConversionPatterns(typeConverter, pattern);
3967+
39623968
mlir::ConversionTarget target{*context};
39633969
target.addLegalDialect<mlir::LLVM::LLVMDialect>();
39643970
// The OpenMP dialect is legal for Operations without regions, for those

0 commit comments

Comments
 (0)