[flang-commits] [flang] [flang] Added alternative inlining code for hlfir.cshift. (PR #129176)
via flang-commits
flang-commits at lists.llvm.org
Thu Feb 27 19:09:15 PST 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-flang-fir-hlfir
Author: Slava Zakharin (vzakhari)
<details>
<summary>Changes</summary>
Flang generates slower code for `CSHIFT(CSHIFT(PTR(:,:,I),sh1,1),sh2,2)`
pattern in facerec than other compilers. The first CSHIFT can be done
as two memcpy's wrapped in a loop for the second dimension.
This does require creating a temporary array, but it seems to be faster,
than the current hlfir.elemental inlining.
I started with modifying the new index computation in
hlfir.elemental inlining: the new arith.select approach does enable
some vectorization in LLVM, but on x86 it is using gathers/scatters
and does not give much speed-up.
I also experimented with LoopBoundSplitPass
and InductiveRangeCheckElimination for a simple (not chained) CSHIFT
case, but I could not adjust them to split the loop with a condition
on the value of the IV into two loops with disjoint iteration spaces.
I thought if I could do it, I would be able to keep the hlfir.elemental
inlining mostly untouched, and then adjust the hlfir.elemental inlining
heuristics for the facerec case.
Since I was not able to make these pass work for me, I added a special
case inlining for CSHIFT(ARRAY,SH,DIM=1) via hlfir.eval_in_mem.
If ARRAY is not statically known to have the contiguous leading
dimension, there is a dynamic check for contiguity, which allows
exposing it to LLVM and enabling the rewrite of the copy loops
into memcpys. This approach is stepping on the toes of LoopVersioning,
but it is helpful in facerec case.
I measured ~6% speed-up on grace, and ~4% on zen4.
---
Patch is 71.10 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/129176.diff
6 Files Affected:
- (modified) flang/include/flang/Optimizer/Builder/HLFIRTools.h (+13)
- (modified) flang/include/flang/Optimizer/Support/DataLayout.h (+6)
- (modified) flang/lib/Optimizer/Builder/HLFIRTools.cpp (+49)
- (modified) flang/lib/Optimizer/HLFIR/Transforms/SimplifyHLFIRIntrinsics.cpp (+331-26)
- (modified) flang/lib/Optimizer/Support/DataLayout.cpp (+17)
- (modified) flang/test/HLFIR/simplify-hlfir-intrinsics-cshift.fir (+247-200)
``````````diff
diff --git a/flang/include/flang/Optimizer/Builder/HLFIRTools.h b/flang/include/flang/Optimizer/Builder/HLFIRTools.h
index 8b1235b50cc6f..1faf451e8b495 100644
--- a/flang/include/flang/Optimizer/Builder/HLFIRTools.h
+++ b/flang/include/flang/Optimizer/Builder/HLFIRTools.h
@@ -517,6 +517,19 @@ Entity loadElementAt(mlir::Location loc, fir::FirOpBuilder &builder,
llvm::SmallVector<mlir::Value, Fortran::common::maxRank>
genExtentsVector(mlir::Location loc, fir::FirOpBuilder &builder, Entity entity);
+/// Generate an hlfir.designate that produces an 1D section
+/// of \p array using \p oneBasedIndices and \p dim:
+/// i = oneBasedIndices
+/// result => array(i(1), ..., i(dim-1), :, i(dim+1), ..., i(n))
+///
+/// The caller provides the pre-computed \p lbounds, \p extents
+/// and \p typeParams of the array.
+Entity gen1DSection(mlir::Location loc, fir::FirOpBuilder &builder,
+ Entity array, int64_t dim,
+ mlir::ArrayRef<mlir::Value> lbounds,
+ mlir::ArrayRef<mlir::Value> extents,
+ mlir::ValueRange oneBasedIndices,
+ mlir::ArrayRef<mlir::Value> typeParams);
} // namespace hlfir
#endif // FORTRAN_OPTIMIZER_BUILDER_HLFIRTOOLS_H
diff --git a/flang/include/flang/Optimizer/Support/DataLayout.h b/flang/include/flang/Optimizer/Support/DataLayout.h
index 957ea99162c5b..64b56f1612589 100644
--- a/flang/include/flang/Optimizer/Support/DataLayout.h
+++ b/flang/include/flang/Optimizer/Support/DataLayout.h
@@ -58,6 +58,12 @@ std::optional<mlir::DataLayout>
getOrSetMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule,
bool allowDefaultLayout = false);
+/// Create mlir::DataLayout from the data layout information on the
+/// mlir::Module. If the DLTI attribute is not set, returns std::nullopt.
+std::optional<mlir::DataLayout> getMLIRDataLayout(mlir::ModuleOp mlirModule);
+std::optional<mlir::DataLayout>
+getMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule);
+
} // namespace fir::support
#endif // FORTRAN_OPTIMIZER_SUPPORT_DATALAYOUT_H
diff --git a/flang/lib/Optimizer/Builder/HLFIRTools.cpp b/flang/lib/Optimizer/Builder/HLFIRTools.cpp
index 8993065c2bb64..f4967ed3852b9 100644
--- a/flang/lib/Optimizer/Builder/HLFIRTools.cpp
+++ b/flang/lib/Optimizer/Builder/HLFIRTools.cpp
@@ -1535,3 +1535,52 @@ hlfir::genExtentsVector(mlir::Location loc, fir::FirOpBuilder &builder,
shape.getDefiningOp()->erase();
return extents;
}
+
+hlfir::Entity hlfir::gen1DSection(mlir::Location loc,
+ fir::FirOpBuilder &builder,
+ hlfir::Entity array, int64_t dim,
+ mlir::ArrayRef<mlir::Value> lbounds,
+ mlir::ArrayRef<mlir::Value> extents,
+ mlir::ValueRange oneBasedIndices,
+ mlir::ArrayRef<mlir::Value> typeParams) {
+ assert(array.isVariable() && "array must be a variable");
+ assert(dim > 0 && dim <= array.getRank() && "invalid dim number");
+ mlir::Value one =
+ builder.createIntegerConstant(loc, builder.getIndexType(), 1);
+ hlfir::DesignateOp::Subscripts subscripts;
+ unsigned indexId = 0;
+ for (int i = 0; i < array.getRank(); ++i) {
+ if (i == dim - 1) {
+ mlir::Value ubound = genUBound(loc, builder, lbounds[i], extents[i], one);
+ subscripts.emplace_back(
+ hlfir::DesignateOp::Triplet{lbounds[i], ubound, one});
+ } else {
+ mlir::Value index =
+ genUBound(loc, builder, lbounds[i], oneBasedIndices[indexId++], one);
+ subscripts.emplace_back(index);
+ }
+ }
+ mlir::Value sectionShape =
+ builder.create<fir::ShapeOp>(loc, extents[dim - 1]);
+
+ // The result type is one of:
+ // !fir.box/class<!fir.array<NxT>>
+ // !fir.box/class<!fir.array<?xT>>
+ //
+ // We could use !fir.ref<!fir.array<NxT>> when the whole dimension's
+ // size is known and it is the leading dimension, but let it be simple
+ // for the time being.
+ auto seqType =
+ mlir::cast<fir::SequenceType>(array.getElementOrSequenceType());
+ int64_t dimExtent = seqType.getShape()[dim - 1];
+ mlir::Type sectionType =
+ fir::SequenceType::get({dimExtent}, seqType.getEleTy());
+ sectionType = fir::wrapInClassOrBoxType(sectionType, array.isPolymorphic());
+
+ auto designate = builder.create<hlfir::DesignateOp>(
+ loc, sectionType, array, /*component=*/"", /*componentShape=*/nullptr,
+ subscripts,
+ /*substring=*/mlir::ValueRange{}, /*complexPartAttr=*/std::nullopt,
+ sectionShape, typeParams);
+ return hlfir::Entity{designate.getResult()};
+}
diff --git a/flang/lib/Optimizer/HLFIR/Transforms/SimplifyHLFIRIntrinsics.cpp b/flang/lib/Optimizer/HLFIR/Transforms/SimplifyHLFIRIntrinsics.cpp
index c1c3839c47e11..b36cd282015e7 100644
--- a/flang/lib/Optimizer/HLFIR/Transforms/SimplifyHLFIRIntrinsics.cpp
+++ b/flang/lib/Optimizer/HLFIR/Transforms/SimplifyHLFIRIntrinsics.cpp
@@ -18,6 +18,7 @@
#include "flang/Optimizer/HLFIR/HLFIRDialect.h"
#include "flang/Optimizer/HLFIR/HLFIROps.h"
#include "flang/Optimizer/HLFIR/Passes.h"
+#include "flang/Optimizer/Support/DataLayout.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/IR/Location.h"
#include "mlir/Pass/Pass.h"
@@ -415,15 +416,13 @@ class SumAsElementalConversion : public mlir::OpRewritePattern<hlfir::SumOp> {
}
};
-class CShiftAsElementalConversion
- : public mlir::OpRewritePattern<hlfir::CShiftOp> {
+class CShiftConversion : public mlir::OpRewritePattern<hlfir::CShiftOp> {
public:
using mlir::OpRewritePattern<hlfir::CShiftOp>::OpRewritePattern;
llvm::LogicalResult
matchAndRewrite(hlfir::CShiftOp cshift,
mlir::PatternRewriter &rewriter) const override {
- using Fortran::common::maxRank;
hlfir::ExprType expr = mlir::dyn_cast<hlfir::ExprType>(cshift.getType());
assert(expr &&
@@ -445,31 +444,78 @@ class CShiftAsElementalConversion
if (dimVal <= 0 || dimVal > arrayRank)
return rewriter.notifyMatchFailure(cshift, "Invalid DIM for CSHIFT");
- mlir::Location loc = cshift.getLoc();
- fir::FirOpBuilder builder{rewriter, cshift.getOperation()};
- mlir::Type elementType = expr.getElementType();
+ // When DIM==1 and the contiguity of the input array is not statically
+ // known, try to exploit the fact that the leading dimension might be
+ // contiguous. We can do this now using hlfir.eval_in_mem with
+ // a dynamic check for the leading dimension contiguity.
+ // Otherwise, convert hlfir.cshift to hlfir.elemental.
+ //
+ // Note that the hlfir.elemental can be inlined into other hlfir.elemental,
+ // while hlfir.eval_in_mem prevents this, and we will end up creating
+ // a temporary array for the result. We may need to come up with
+ // a more sophisticated logic for picking the most efficient
+ // representation.
hlfir::Entity array = hlfir::Entity{cshift.getArray()};
- mlir::Value arrayShape = hlfir::genShape(loc, builder, array);
- llvm::SmallVector<mlir::Value> arrayExtents =
- hlfir::getExplicitExtentsFromShape(arrayShape, builder);
- llvm::SmallVector<mlir::Value, 1> typeParams;
- hlfir::genLengthParameters(loc, builder, array, typeParams);
+ mlir::Type elementType = array.getFortranElementType();
+ if (dimVal == 1 && fir::isa_trivial(elementType) &&
+ !array.isSimplyContiguous())
+ rewriter.replaceOp(cshift, genInMemCShift(rewriter, cshift, dimVal));
+ else
+ rewriter.replaceOp(cshift, genElementalCShift(rewriter, cshift, dimVal));
+ return mlir::success();
+ }
+
+private:
+ /// Generate MODULO(\p shiftVal, \p extent).
+ static mlir::Value normalizeShiftValue(mlir::Location loc,
+ fir::FirOpBuilder &builder,
+ mlir::Value shiftVal,
+ mlir::Value extent,
+ mlir::Type calcType) {
+ shiftVal = builder.createConvert(loc, calcType, shiftVal);
+ extent = builder.createConvert(loc, calcType, extent);
+ shiftVal = fir::IntrinsicLibrary{builder, loc}.genModulo(
+ calcType, {shiftVal, extent});
+ return builder.createConvert(loc, calcType, shiftVal);
+ }
+
+ /// Convert \p cshift into an hlfir.elemental using
+ /// the pre-computed constant \p dimVal.
+ static mlir::Operation *genElementalCShift(mlir::PatternRewriter &rewriter,
+ hlfir::CShiftOp cshift,
+ int64_t dimVal) {
+ using Fortran::common::maxRank;
hlfir::Entity shift = hlfir::Entity{cshift.getShift()};
+ hlfir::Entity array = hlfir::Entity{cshift.getArray()};
+
+ mlir::Location loc = cshift.getLoc();
+ fir::FirOpBuilder builder{rewriter, cshift.getOperation()};
// The new index computation involves MODULO, which is not implemented
// for IndexType, so use I64 instead.
mlir::Type calcType = builder.getI64Type();
+ // All the indices arithmetic used below does not overflow
+ // signed and unsigned I64.
+ builder.setIntegerOverflowFlags(mlir::arith::IntegerOverflowFlags::nsw |
+ mlir::arith::IntegerOverflowFlags::nuw);
- mlir::Value one = builder.createIntegerConstant(loc, calcType, 1);
+ mlir::Value arrayShape = hlfir::genShape(loc, builder, array);
+ llvm::SmallVector<mlir::Value, maxRank> arrayExtents =
+ hlfir::getExplicitExtentsFromShape(arrayShape, builder);
+ llvm::SmallVector<mlir::Value, 1> typeParams;
+ hlfir::genLengthParameters(loc, builder, array, typeParams);
+ mlir::Value shiftDimExtent =
+ builder.createConvert(loc, calcType, arrayExtents[dimVal - 1]);
mlir::Value shiftVal;
if (shift.isScalar()) {
shiftVal = hlfir::loadTrivialScalar(loc, builder, shift);
- shiftVal = builder.createConvert(loc, calcType, shiftVal);
+ shiftVal =
+ normalizeShiftValue(loc, builder, shiftVal, shiftDimExtent, calcType);
}
auto genKernel = [&](mlir::Location loc, fir::FirOpBuilder &builder,
mlir::ValueRange inputIndices) -> hlfir::Entity {
llvm::SmallVector<mlir::Value, maxRank> indices{inputIndices};
- if (!shift.isScalar()) {
+ if (!shiftVal) {
// When the array is not a vector, section
// (s(1), s(2), ..., s(dim-1), :, s(dim+1), ..., s(n)
// of the result has a value equal to:
@@ -482,35 +528,294 @@ class CShiftAsElementalConversion
hlfir::Entity shiftElement =
hlfir::getElementAt(loc, builder, shift, shiftIndices);
shiftVal = hlfir::loadTrivialScalar(loc, builder, shiftElement);
- shiftVal = builder.createConvert(loc, calcType, shiftVal);
+ shiftVal = normalizeShiftValue(loc, builder, shiftVal, shiftDimExtent,
+ calcType);
}
// Element i of the result (1-based) is element
- // 'MODULO(i + SH - 1, SIZE(ARRAY)) + 1' (1-based) of the original
+ // 'MODULO(i + SH - 1, SIZE(ARRAY,DIM)) + 1' (1-based) of the original
// ARRAY (or its section, when ARRAY is not a vector).
+
+ // Compute the index into the original array using the normalized
+ // shift value, which satisfies (SH >= 0 && SH < SIZE(ARRAY,DIM)):
+ // newIndex =
+ // i + ((i <= SIZE(ARRAY,DIM) - SH) ? SH : SH - SIZE(ARRAY,DIM))
+ //
+ // Such index computation allows for further loop vectorization
+ // in LLVM.
+ mlir::Value wrapBound =
+ builder.create<mlir::arith::SubIOp>(loc, shiftDimExtent, shiftVal);
+ mlir::Value adjustedShiftVal =
+ builder.create<mlir::arith::SubIOp>(loc, shiftVal, shiftDimExtent);
mlir::Value index =
builder.createConvert(loc, calcType, inputIndices[dimVal - 1]);
- mlir::Value extent = arrayExtents[dimVal - 1];
+ mlir::Value wrapCheck = builder.create<mlir::arith::CmpIOp>(
+ loc, mlir::arith::CmpIPredicate::sle, index, wrapBound);
+ mlir::Value actualShift = builder.create<mlir::arith::SelectOp>(
+ loc, wrapCheck, shiftVal, adjustedShiftVal);
mlir::Value newIndex =
- builder.create<mlir::arith::AddIOp>(loc, index, shiftVal);
- newIndex = builder.create<mlir::arith::SubIOp>(loc, newIndex, one);
- newIndex = fir::IntrinsicLibrary{builder, loc}.genModulo(
- calcType, {newIndex, builder.createConvert(loc, calcType, extent)});
- newIndex = builder.create<mlir::arith::AddIOp>(loc, newIndex, one);
+ builder.create<mlir::arith::AddIOp>(loc, index, actualShift);
newIndex = builder.createConvert(loc, builder.getIndexType(), newIndex);
-
indices[dimVal - 1] = newIndex;
hlfir::Entity element = hlfir::getElementAt(loc, builder, array, indices);
return hlfir::loadTrivialScalar(loc, builder, element);
};
+ mlir::Type elementType = array.getFortranElementType();
hlfir::ElementalOp elementalOp = hlfir::genElementalOp(
loc, builder, elementType, arrayShape, typeParams, genKernel,
/*isUnordered=*/true,
array.isPolymorphic() ? static_cast<mlir::Value>(array) : nullptr,
cshift.getResult().getType());
- rewriter.replaceOp(cshift, elementalOp);
- return mlir::success();
+ return elementalOp.getOperation();
+ }
+
+ /// Convert \p cshift into an hlfir.eval_in_mem using the pre-computed
+ /// constant \p dimVal.
+ /// The converted code looks like this:
+ /// do i=1,SH
+ /// result(i + (SIZE(ARRAY,DIM) - SH)) = array(i)
+ /// end
+ /// do i=1,SIZE(ARRAY,DIM) - SH
+ /// result(i) = array(i + SH)
+ /// end
+ ///
+ /// When \p dimVal is 1, we generate the same code twice
+ /// under a dynamic check for the contiguity of the leading
+ /// dimension. In the code corresponding to the contiguous
+ /// leading dimension, the shift dimension is represented
+ /// as a contiguous slice of the original array.
+ /// This allows recognizing the above two loops as memcpy
+ /// loop idioms in LLVM.
+ static mlir::Operation *genInMemCShift(mlir::PatternRewriter &rewriter,
+ hlfir::CShiftOp cshift,
+ int64_t dimVal) {
+ using Fortran::common::maxRank;
+ hlfir::Entity shift = hlfir::Entity{cshift.getShift()};
+ hlfir::Entity array = hlfir::Entity{cshift.getArray()};
+ assert(array.isVariable() && "array must be a variable");
+ assert(!array.isPolymorphic() &&
+ "genInMemCShift does not support polymorphic types");
+ mlir::Location loc = cshift.getLoc();
+ fir::FirOpBuilder builder{rewriter, cshift.getOperation()};
+ // The new index computation involves MODULO, which is not implemented
+ // for IndexType, so use I64 instead.
+ mlir::Type calcType = builder.getI64Type();
+ // All the indices arithmetic used below does not overflow
+ // signed and unsigned I64.
+ builder.setIntegerOverflowFlags(mlir::arith::IntegerOverflowFlags::nsw |
+ mlir::arith::IntegerOverflowFlags::nuw);
+
+ mlir::Value arrayShape = hlfir::genShape(loc, builder, array);
+ llvm::SmallVector<mlir::Value, maxRank> arrayExtents =
+ hlfir::getExplicitExtentsFromShape(arrayShape, builder);
+ llvm::SmallVector<mlir::Value, 1> typeParams;
+ hlfir::genLengthParameters(loc, builder, array, typeParams);
+ mlir::Value shiftDimExtent =
+ builder.createConvert(loc, calcType, arrayExtents[dimVal - 1]);
+ mlir::Value shiftVal;
+ if (shift.isScalar()) {
+ shiftVal = hlfir::loadTrivialScalar(loc, builder, shift);
+ shiftVal =
+ normalizeShiftValue(loc, builder, shiftVal, shiftDimExtent, calcType);
+ }
+
+ hlfir::EvaluateInMemoryOp evalOp =
+ builder.create<hlfir::EvaluateInMemoryOp>(
+ loc, mlir::cast<hlfir::ExprType>(cshift.getType()), arrayShape);
+ builder.setInsertionPointToStart(&evalOp.getBody().front());
+
+ mlir::Value resultArray = evalOp.getMemory();
+ mlir::Type arrayType = fir::dyn_cast_ptrEleTy(resultArray.getType());
+ resultArray = builder.createBox(loc, fir::BoxType::get(arrayType),
+ resultArray, arrayShape, /*slice=*/nullptr,
+ typeParams, /*tdesc=*/nullptr);
+
+ // This is a generator of the dimension shift code.
+ // The code is inserted inside a loop nest over the other dimensions
+ // (if any). If exposeContiguity is true, the array's section
+ // array(s(1), ..., s(dim-1), :, s(dim+1), ..., s(n)) is represented
+ // as a contiguous 1D array.
+ // shiftVal is the normalized shift value that satisfies (SH >= 0 && SH <
+ // SIZE(ARRAY,DIM)).
+ //
+ auto genDimensionShift = [&](mlir::Location loc, fir::FirOpBuilder &builder,
+ mlir::Value shiftVal, bool exposeContiguity,
+ mlir::ValueRange oneBasedIndices)
+ -> llvm::SmallVector<mlir::Value, 0> {
+ // Create a vector of indices (s(1), ..., s(dim-1), nullptr, s(dim+1),
+ // ..., s(n)) so that we can update the dimVal index as needed.
+ llvm::SmallVector<mlir::Value, maxRank> srcIndices(
+ oneBasedIndices.begin(), oneBasedIndices.begin() + (dimVal - 1));
+ srcIndices.push_back(nullptr);
+ srcIndices.append(oneBasedIndices.begin() + (dimVal - 1),
+ oneBasedIndices.end());
+ llvm::SmallVector<mlir::Value, maxRank> dstIndices(srcIndices);
+
+ hlfir::Entity srcArray = array;
+ if (exposeContiguity && mlir::isa<fir::BaseBoxType>(srcArray.getType())) {
+ assert(dimVal == 1 && "can expose contiguity only for dim 1");
+ llvm::SmallVector<mlir::Value, maxRank> arrayLbounds =
+ hlfir::genLowerbounds(loc, builder, arrayShape, array.getRank());
+ hlfir::Entity section =
+ hlfir::gen1DSection(loc, builder, srcArray, dimVal, arrayLbounds,
+ arrayExtents, oneBasedIndices, typeParams);
+ mlir::Value addr = hlfir::genVariableRawAddress(loc, builder, section);
+ mlir::Value shape = hlfir::genShape(loc, builder, section);
+ mlir::Type boxType = fir::wrapInClassOrBoxType(
+ hlfir::getFortranElementOrSequenceType(section.getType()),
+ section.isPolymorphic());
+ srcArray = hlfir::Entity{
+ builder.createBox(loc, boxType, addr, shape, /*slice=*/nullptr,
+ /*lengths=*/{}, /*tdesc=*/nullptr)};
+ // When shifting the dimension as a 1D section of the original
+ // array, we only need one index for addressing.
+ srcIndices.resize(1);
+ }
+
+ // Copy first portion of the array:
+ // do i=1,SH
+ // result(i + (SIZE(ARRAY,DIM) - SH)) = array(i)
+ // end
+ auto genAssign1 = [&](mlir::Location loc, fir::FirOpBuilder &builder,
+ mlir::ValueRange index,
+ mlir::ValueRange reductionArgs)
+ -> llvm::SmallVector<mlir::Value, 0> {
+ assert(index.size() == 1 && "expected single loop");
+ mlir::Value srcIndex = builder.createConvert(loc, calcType, index[0]);
+ srcIndices[dimVal - 1] = srcIndex;
+ hlfir::Entity srcElementValue =
+ hlfir::loadElementAt(loc, builder, srcArray, srcIndices);
+ mlir::Value dstIndex = builder.create<mlir::arith::AddIOp>(
+ loc, srcIndex,
+ builder.create<mlir::arith::SubIOp>(loc, shiftDimExtent, shiftVal));
+ dstIndices[dimVal - 1] = dstIndex;
+ hlfir::Entity dstElement = hlfir::getElementAt(
+ loc, builder, hlfir::Entity{resultArray}, dstIndices);
+ builder.create<hlfir::AssignOp>(loc, srcElementValue, dstElement);
+ return {};
+ };
+
+ // Generate the first loop.
+ hlfir::genLoopNestWithReductions(loc, builder, {shiftVal},
+ /*reductionInits=*/{}, genAssign1,
+ /*isUnordered=*/true);
+
+ // Copy second portion of the array:
+ // do i=1,SIZE(ARRAY,DIM)-SH
+ // result(i) = array(i + SH)
+ // end
+ auto genAssign2 = [&](mlir::Location loc, fir::FirOpBuilder...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/129176
More information about the flang-commits
mailing list