Skip to content

Commit e234e6e

Browse files
vzakharivzakhari
committed
Use BoxEleSizeOp, and moved a misplaced isSimplyContiguous() check.
1 parent c8c2b57 commit e234e6e

File tree

4 files changed

+14
-54
lines changed

4 files changed

+14
-54
lines changed

‎flang/include/flang/Optimizer/Support/DataLayout.h‎

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -58,12 +58,6 @@ std::optional<mlir::DataLayout>
5858
getOrSetMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule,
5959
bool allowDefaultLayout = false);
6060

61-
/// Create mlir::DataLayout from the data layout information on the
62-
/// mlir::Module. If the DLTI attribute is not set, returns std::nullopt.
63-
std::optional<mlir::DataLayout> getMLIRDataLayout(mlir::ModuleOp mlirModule);
64-
std::optional<mlir::DataLayout>
65-
getMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule);
66-
6761
} // namespace fir::support
6862

6963
#endif // FORTRAN_OPTIMIZER_SUPPORT_DATALAYOUT_H

‎flang/lib/Optimizer/HLFIR/Transforms/SimplifyHLFIRIntrinsics.cpp‎

Lines changed: 11 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,6 @@
1818
#include "flang/Optimizer/HLFIR/HLFIRDialect.h"
1919
#include "flang/Optimizer/HLFIR/HLFIROps.h"
2020
#include "flang/Optimizer/HLFIR/Passes.h"
21-
#include "flang/Optimizer/Support/DataLayout.h"
2221
#include "mlir/Dialect/Arith/IR/Arith.h"
2322
#include "mlir/IR/Location.h"
2423
#include "mlir/Pass/Pass.h"
@@ -457,8 +456,7 @@ class CShiftConversion : public mlir::OpRewritePattern<hlfir::CShiftOp> {
457456
// representation.
458457
hlfir::Entity array = hlfir::Entity{cshift.getArray()};
459458
mlir::Type elementType = array.getFortranElementType();
460-
if (dimVal == 1 && fir::isa_trivial(elementType) &&
461-
!array.isSimplyContiguous())
459+
if (dimVal == 1 && fir::isa_trivial(elementType))
462460
rewriter.replaceOp(cshift, genInMemCShift(rewriter, cshift, dimVal));
463461
else
464462
rewriter.replaceOp(cshift, genElementalCShift(rewriter, cshift, dimVal));
@@ -759,30 +757,18 @@ class CShiftConversion : public mlir::OpRewritePattern<hlfir::CShiftOp> {
759757
mlir::Value elemSize;
760758
mlir::Value stride;
761759
mlir::Type elementType = array.getFortranElementType();
762-
if (dimVal == 1 && mlir::isa<fir::BaseBoxType>(array.getType()) &&
763-
fir::isa_trivial(elementType)) {
764-
mlir::ModuleOp module = cshift->getParentOfType<mlir::ModuleOp>();
765-
std::optional<mlir::DataLayout> dl =
766-
fir::support::getMLIRDataLayout(module);
767-
if (dl) {
768-
fir::KindMapping kindMap = fir::getKindMapping(module);
769-
auto [size, align] = fir::getTypeSizeAndAlignmentOrCrash(
770-
loc, elementType, *dl, kindMap);
771-
size = llvm::alignTo(size, align);
772-
if (size) {
773-
mlir::Type indexType = builder.getIndexType();
774-
elemSize = builder.createIntegerConstant(loc, indexType, size);
775-
776-
mlir::Value dimIdx =
777-
builder.createIntegerConstant(loc, indexType, dimVal - 1);
778-
auto boxDim = builder.create<fir::BoxDimsOp>(
779-
loc, indexType, indexType, indexType, array.getBase(), dimIdx);
780-
stride = boxDim.getByteStride();
781-
}
782-
}
760+
if (dimVal == 1 && mlir::isa<fir::BaseBoxType>(array.getType())) {
761+
mlir::Type indexType = builder.getIndexType();
762+
elemSize =
763+
builder.create<fir::BoxEleSizeOp>(loc, indexType, array.getBase());
764+
mlir::Value dimIdx =
765+
builder.createIntegerConstant(loc, indexType, dimVal - 1);
766+
auto boxDim = builder.create<fir::BoxDimsOp>(
767+
loc, indexType, indexType, indexType, array.getBase(), dimIdx);
768+
stride = boxDim.getByteStride();
783769
}
784770

785-
if (!elemSize || !stride) {
771+
if (array.isSimplyContiguous() || !elemSize || !stride) {
786772
genDimensionShift(loc, builder, shiftVal, /*exposeContiguity=*/false,
787773
oneBasedIndices);
788774
return {};

‎flang/lib/Optimizer/Support/DataLayout.cpp‎

Lines changed: 0 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -50,13 +50,6 @@ static void setDataLayoutFromAttributes(ModOpTy mlirModule,
5050
fir::support::setMLIRDataLayout(mlirModule, llvmDataLayout);
5151
}
5252

53-
template <typename ModOpTy>
54-
static std::optional<mlir::DataLayout> getDataLayout(ModOpTy mlirModule) {
55-
if (!mlirModule.getDataLayoutSpec())
56-
return std::nullopt;
57-
return mlir::DataLayout(mlirModule);
58-
}
59-
6053
template <typename ModOpTy>
6154
static std::optional<mlir::DataLayout>
6255
getOrSetDataLayout(ModOpTy mlirModule, bool allowDefaultLayout) {
@@ -102,13 +95,3 @@ fir::support::getOrSetMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule,
10295
bool allowDefaultLayout) {
10396
return getOrSetDataLayout(mlirModule, allowDefaultLayout);
10497
}
105-
106-
std::optional<mlir::DataLayout>
107-
fir::support::getMLIRDataLayout(mlir::ModuleOp mlirModule) {
108-
return getDataLayout(mlirModule);
109-
}
110-
111-
std::optional<mlir::DataLayout>
112-
fir::support::getMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule) {
113-
return getDataLayout(mlirModule);
114-
}

‎flang/test/HLFIR/simplify-hlfir-intrinsics-cshift.fir‎

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,6 @@
11
// Test hlfir.cshift simplification to hlfir.elemental and hlfir.eval_in_mem:
22
// RUN: fir-opt --simplify-hlfir-intrinsics %s | FileCheck %s
33

4-
module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} {
5-
64
func.func @cshift_vector(%arg0: !fir.box<!fir.array<?xi32>>, %arg1: !fir.ref<i32>) -> !hlfir.expr<?xi32>{
75
%res = hlfir.cshift %arg0 %arg1 : (!fir.box<!fir.array<?xi32>>, !fir.ref<i32>) -> !hlfir.expr<?xi32>
86
return %res : !hlfir.expr<?xi32>
@@ -11,7 +9,6 @@ func.func @cshift_vector(%arg0: !fir.box<!fir.array<?xi32>>, %arg1: !fir.ref<i32
119
// CHECK-SAME: %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !fir.box<!fir.array<?xi32>>,
1210
// CHECK-SAME: %[[VAL_1:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !fir.ref<i32>) -> !hlfir.expr<?xi32> {
1311
// CHECK: %[[VAL_2:.*]] = arith.constant 1 : index
14-
// CHECK: %[[VAL_3:.*]] = arith.constant 4 : index
1512
// CHECK: %[[VAL_4:.*]] = arith.constant 0 : i64
1613
// CHECK: %[[VAL_5:.*]] = arith.constant 0 : index
1714
// CHECK: %[[VAL_6:.*]]:3 = fir.box_dims %[[VAL_0]], %[[VAL_5]] : (!fir.box<!fir.array<?xi32>>, index) -> (index, index, index)
@@ -29,8 +26,10 @@ func.func @cshift_vector(%arg0: !fir.box<!fir.array<?xi32>>, %arg1: !fir.ref<i32
2926
// CHECK: %[[VAL_18:.*]] = hlfir.eval_in_mem shape %[[VAL_7]] : (!fir.shape<1>) -> !hlfir.expr<?xi32> {
3027
// CHECK: ^bb0(%[[VAL_19:.*]]: !fir.ref<!fir.array<?xi32>>):
3128
// CHECK: %[[VAL_20:.*]] = fir.embox %[[VAL_19]](%[[VAL_7]]) : (!fir.ref<!fir.array<?xi32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xi32>>
29+
// CHECK: %[[VAL_3:.*]] = fir.box_elesize %[[VAL_0]] : (!fir.box<!fir.array<?xi32>>) -> index
3230
// CHECK: %[[VAL_21:.*]]:3 = fir.box_dims %[[VAL_0]], %[[VAL_5]] : (!fir.box<!fir.array<?xi32>>, index) -> (index, index, index)
33-
// CHECK: %[[VAL_22:.*]] = arith.cmpi eq, %[[VAL_21]]#2, %[[VAL_3]] : index
31+
32+
// CHECK: %[[VAL_22:.*]] = arith.cmpi eq, %[[VAL_3]], %[[VAL_21]]#2 : index
3433
// CHECK: fir.if %[[VAL_22]] {
3534
// CHECK: %[[VAL_23:.*]] = fir.shape %[[VAL_6]]#1 : (index) -> !fir.shape<1>
3635
// CHECK: %[[VAL_24:.*]] = hlfir.designate %[[VAL_0]] (%[[VAL_2]]:%[[VAL_6]]#1:%[[VAL_2]]) shape %[[VAL_23]] : (!fir.box<!fir.array<?xi32>>, index, index, index, !fir.shape<1>) -> !fir.box<!fir.array<?xi32>>
@@ -321,5 +320,3 @@ func.func @cshift_vector_assumed_dim_1(%arg0: !fir.box<!fir.array<?xi32>>, %arg1
321320
}
322321
// CHECK-LABEL: func.func @cshift_vector_assumed_dim_1(
323322
// CHECK-NOT: hlfir.cshift
324-
325-
} // end module

0 commit comments

Comments
 (0)