aboutsummaryrefslogtreecommitdiff
path: root/flang/lib/Optimizer/CodeGen
diff options
context:
space:
mode:
Diffstat (limited to 'flang/lib/Optimizer/CodeGen')
-rw-r--r--flang/lib/Optimizer/CodeGen/CodeGen.cpp41
-rw-r--r--flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp7
-rw-r--r--flang/lib/Optimizer/CodeGen/TargetRewrite.cpp24
3 files changed, 64 insertions, 8 deletions
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 70bb43a2..478ab15 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -749,6 +749,44 @@ struct VolatileCastOpConversion
}
};
+/// Lower `fir.assumed_size_extent` to constant -1 of index type.
+struct AssumedSizeExtentOpConversion
+ : public fir::FIROpConversion<fir::AssumedSizeExtentOp> {
+ using FIROpConversion::FIROpConversion;
+
+ llvm::LogicalResult
+ matchAndRewrite(fir::AssumedSizeExtentOp op, OpAdaptor,
+ mlir::ConversionPatternRewriter &rewriter) const override {
+ mlir::Location loc = op.getLoc();
+ mlir::Type ity = lowerTy().indexType();
+ auto cst = fir::genConstantIndex(loc, ity, rewriter, -1);
+ rewriter.replaceOp(op, cst.getResult());
+ return mlir::success();
+ }
+};
+
+/// Lower `fir.is_assumed_size_extent` to integer equality with -1.
+struct IsAssumedSizeExtentOpConversion
+ : public fir::FIROpConversion<fir::IsAssumedSizeExtentOp> {
+ using FIROpConversion::FIROpConversion;
+
+ llvm::LogicalResult
+ matchAndRewrite(fir::IsAssumedSizeExtentOp op, OpAdaptor adaptor,
+ mlir::ConversionPatternRewriter &rewriter) const override {
+ mlir::Location loc = op.getLoc();
+ mlir::Value val = adaptor.getVal();
+ mlir::Type valTy = val.getType();
+ // Create constant -1 of the operand type.
+ auto negOneAttr = rewriter.getIntegerAttr(valTy, -1);
+ auto negOne =
+ mlir::LLVM::ConstantOp::create(rewriter, loc, valTy, negOneAttr);
+ auto cmp = mlir::LLVM::ICmpOp::create(
+ rewriter, loc, mlir::LLVM::ICmpPredicate::eq, val, negOne);
+ rewriter.replaceOp(op, cmp.getResult());
+ return mlir::success();
+ }
+};
+
/// convert value of from-type to value of to-type
struct ConvertOpConversion : public fir::FIROpConversion<fir::ConvertOp> {
using FIROpConversion::FIROpConversion;
@@ -1113,7 +1151,7 @@ struct AllocMemOpConversion : public fir::FIROpConversion<fir::AllocMemOp> {
mlir::Value size = genTypeSizeInBytes(loc, ity, rewriter, llvmObjectTy);
if (auto scaleSize =
fir::genAllocationScaleSize(loc, heap.getInType(), ity, rewriter))
- size = rewriter.create<mlir::LLVM::MulOp>(loc, ity, size, scaleSize);
+ size = mlir::LLVM::MulOp::create(rewriter, loc, ity, size, scaleSize);
for (mlir::Value opnd : adaptor.getOperands())
size = mlir::LLVM::MulOp::create(rewriter, loc, ity, size,
integerCast(loc, rewriter, ity, opnd));
@@ -4360,6 +4398,7 @@ void fir::populateFIRToLLVMConversionPatterns(
AllocaOpConversion, AllocMemOpConversion, BoxAddrOpConversion,
BoxCharLenOpConversion, BoxDimsOpConversion, BoxEleSizeOpConversion,
BoxIsAllocOpConversion, BoxIsArrayOpConversion, BoxIsPtrOpConversion,
+ AssumedSizeExtentOpConversion, IsAssumedSizeExtentOpConversion,
BoxOffsetOpConversion, BoxProcHostOpConversion, BoxRankOpConversion,
BoxTypeCodeOpConversion, BoxTypeDescOpConversion, CallOpConversion,
CmpcOpConversion, VolatileCastOpConversion, ConvertOpConversion,
diff --git a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
index 381b2a2..f74d635 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
@@ -242,10 +242,11 @@ struct TargetAllocMemOpConversion
loc, llvmObjectTy, ity, rewriter, lowerTy().getDataLayout());
if (auto scaleSize = fir::genAllocationScaleSize(
loc, allocmemOp.getInType(), ity, rewriter))
- size = rewriter.create<mlir::LLVM::MulOp>(loc, ity, size, scaleSize);
+ size = mlir::LLVM::MulOp::create(rewriter, loc, ity, size, scaleSize);
for (mlir::Value opnd : adaptor.getOperands().drop_front())
- size = rewriter.create<mlir::LLVM::MulOp>(
- loc, ity, size, integerCast(lowerTy(), loc, rewriter, ity, opnd));
+ size = mlir::LLVM::MulOp::create(
+ rewriter, loc, ity, size,
+ integerCast(lowerTy(), loc, rewriter, ity, opnd));
auto mallocTyWidth = lowerTy().getIndexTypeBitwidth();
auto mallocTy =
mlir::IntegerType::get(rewriter.getContext(), mallocTyWidth);
diff --git a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
index ac285b5..0776346 100644
--- a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
+++ b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
@@ -872,6 +872,14 @@ public:
}
}
+ // Count the number of arguments that have to stay in place at the end of
+ // the argument list.
+ unsigned trailingArgs = 0;
+ if constexpr (std::is_same_v<FuncOpTy, mlir::gpu::GPUFuncOp>) {
+ trailingArgs =
+ func.getNumWorkgroupAttributions() + func.getNumPrivateAttributions();
+ }
+
// Convert return value(s)
for (auto ty : funcTy.getResults())
llvm::TypeSwitch<mlir::Type>(ty)
@@ -981,6 +989,16 @@ public:
}
}
+ // Add the argument at the end if the number of trailing arguments is 0,
+ // otherwise insert the argument at the appropriate index.
+ auto addOrInsertArgument = [&](mlir::Type ty, mlir::Location loc) {
+ unsigned inputIndex = func.front().getArguments().size() - trailingArgs;
+ auto newArg = trailingArgs == 0
+ ? func.front().addArgument(ty, loc)
+ : func.front().insertArgument(inputIndex, ty, loc);
+ return newArg;
+ };
+
if (!func.empty()) {
// If the function has a body, then apply the fixups to the arguments and
// return ops as required. These fixups are done in place.
@@ -1117,8 +1135,7 @@ public:
// original arguments. (Boxchar arguments.)
auto newBufArg =
func.front().insertArgument(fixup.index, fixupType, loc);
- auto newLenArg =
- func.front().addArgument(trailingTys[fixup.second], loc);
+ auto newLenArg = addOrInsertArgument(trailingTys[fixup.second], loc);
auto boxTy = oldArgTys[fixup.index - offset];
rewriter->setInsertionPointToStart(&func.front());
auto box = fir::EmboxCharOp::create(*rewriter, loc, boxTy, newBufArg,
@@ -1133,8 +1150,7 @@ public:
// appended after all the original arguments.
auto newProcPointerArg =
func.front().insertArgument(fixup.index, fixupType, loc);
- auto newLenArg =
- func.front().addArgument(trailingTys[fixup.second], loc);
+ auto newLenArg = addOrInsertArgument(trailingTys[fixup.second], loc);
auto tupleType = oldArgTys[fixup.index - offset];
rewriter->setInsertionPointToStart(&func.front());
fir::FirOpBuilder builder(*rewriter, getModule());
generated by cgit v1.1 at 2025-10-23 10:02:35 +0000