aboutsummaryrefslogtreecommitdiff
path: root/flang/lib/Optimizer/Builder
diff options
context:
space:
mode:
Diffstat (limited to 'flang/lib/Optimizer/Builder')
-rw-r--r--flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp268
-rw-r--r--flang/lib/Optimizer/Builder/CUFCommon.cpp64
-rw-r--r--flang/lib/Optimizer/Builder/FIRBuilder.cpp37
-rw-r--r--flang/lib/Optimizer/Builder/HLFIRTools.cpp70
-rw-r--r--flang/lib/Optimizer/Builder/IntrinsicCall.cpp135
-rw-r--r--flang/lib/Optimizer/Builder/Runtime/Character.cpp23
-rw-r--r--flang/lib/Optimizer/Builder/Runtime/Intrinsics.cpp40
-rw-r--r--flang/lib/Optimizer/Builder/Runtime/Reduction.cpp2
8 files changed, 538 insertions, 101 deletions
diff --git a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
index 4e276a7..3c86a9d 100644
--- a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
@@ -17,6 +17,8 @@
#include "flang/Evaluate/common.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/MutableBox.h"
+#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
+#include "flang/Optimizer/HLFIR/HLFIROps.h"
#include "mlir/Dialect/Index/IR/IndexOps.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
@@ -51,6 +53,8 @@ static const char __ldlu_r2x2[] = "__ldlu_r2x2_";
static const char __ldlu_r4x4[] = "__ldlu_r4x4_";
static const char __ldlu_r8x2[] = "__ldlu_r8x2_";
+static constexpr unsigned kTMAAlignment = 16;
+
// CUDA specific intrinsic handlers.
static constexpr IntrinsicHandler cudaHandlers[]{
{"__ldca_i4x4",
@@ -195,7 +199,7 @@ static constexpr IntrinsicHandler cudaHandlers[]{
false},
{"atomicadd_r4x4",
static_cast<CUDAIntrinsicLibrary::ExtendedGenerator>(
- &CI::genAtomicAddVector<4>),
+ &CI::genAtomicAddVector4x4),
{{{"a", asAddr}, {"v", asAddr}}},
false},
{"atomicaddd",
@@ -368,6 +372,16 @@ static constexpr IntrinsicHandler cudaHandlers[]{
&CI::genNVVMTime<mlir::NVVM::Clock64Op>),
{},
/*isElemental=*/false},
+ {"cluster_block_index",
+ static_cast<CUDAIntrinsicLibrary::ElementalGenerator>(
+ &CI::genClusterBlockIndex),
+ {},
+ /*isElemental=*/false},
+ {"cluster_dim_blocks",
+ static_cast<CUDAIntrinsicLibrary::ElementalGenerator>(
+ &CI::genClusterDimBlocks),
+ {},
+ /*isElemental=*/false},
{"fence_proxy_async",
static_cast<CUDAIntrinsicLibrary::SubroutineGenerator>(
&CI::genFenceProxyAsync),
@@ -457,6 +471,10 @@ static constexpr IntrinsicHandler cudaHandlers[]{
static_cast<CUDAIntrinsicLibrary::SubroutineGenerator>(&CI::genSyncWarp),
{},
/*isElemental=*/false},
+ {"this_cluster",
+ static_cast<CUDAIntrinsicLibrary::ElementalGenerator>(&CI::genThisCluster),
+ {},
+ /*isElemental=*/false},
{"this_grid",
static_cast<CUDAIntrinsicLibrary::ElementalGenerator>(&CI::genThisGrid),
{},
@@ -472,17 +490,17 @@ static constexpr IntrinsicHandler cudaHandlers[]{
/*isElemental=*/false},
{"threadfence",
static_cast<CUDAIntrinsicLibrary::SubroutineGenerator>(
- &CI::genThreadFence),
+ &CI::genThreadFence<mlir::NVVM::MemScopeKind::GPU>),
{},
/*isElemental=*/false},
{"threadfence_block",
static_cast<CUDAIntrinsicLibrary::SubroutineGenerator>(
- &CI::genThreadFenceBlock),
+ &CI::genThreadFence<mlir::NVVM::MemScopeKind::CTA>),
{},
/*isElemental=*/false},
{"threadfence_system",
static_cast<CUDAIntrinsicLibrary::SubroutineGenerator>(
- &CI::genThreadFenceSystem),
+ &CI::genThreadFence<mlir::NVVM::MemScopeKind::SYS>),
{},
/*isElemental=*/false},
{"tma_bulk_commit_group",
@@ -744,6 +762,56 @@ fir::ExtendedValue CUDAIntrinsicLibrary::genAtomicAddVector(
return fir::ArrayBoxValue(res, {ext});
}
+// ATOMICADDVECTOR4x4
+fir::ExtendedValue CUDAIntrinsicLibrary::genAtomicAddVector4x4(
+ mlir::Type resultType, llvm::ArrayRef<fir::ExtendedValue> args) {
+ assert(args.size() == 2);
+ mlir::Value a = fir::getBase(args[0]);
+ if (mlir::isa<fir::BaseBoxType>(a.getType()))
+ a = fir::BoxAddrOp::create(builder, loc, a);
+
+ const unsigned extent = 4;
+ auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(builder.getContext());
+ mlir::Value ptr = builder.createConvert(loc, llvmPtrTy, a);
+ mlir::Type f32Ty = builder.getF32Type();
+ mlir::Type idxTy = builder.getIndexType();
+ mlir::Type refTy = fir::ReferenceType::get(f32Ty);
+ llvm::SmallVector<mlir::Value> values;
+ for (unsigned i = 0; i < extent; ++i) {
+ mlir::Value pos = builder.createIntegerConstant(loc, idxTy, i);
+ mlir::Value coord = fir::CoordinateOp::create(builder, loc, refTy,
+ fir::getBase(args[1]), pos);
+ mlir::Value value = fir::LoadOp::create(builder, loc, coord);
+ values.push_back(value);
+ }
+
+ auto inlinePtx = mlir::NVVM::InlinePtxOp::create(
+ builder, loc, {f32Ty, f32Ty, f32Ty, f32Ty},
+ {ptr, values[0], values[1], values[2], values[3]}, {},
+ "atom.add.v4.f32 {%0, %1, %2, %3}, [%4], {%5, %6, %7, %8};", {});
+
+ llvm::SmallVector<mlir::Value> results;
+ results.push_back(inlinePtx.getResult(0));
+ results.push_back(inlinePtx.getResult(1));
+ results.push_back(inlinePtx.getResult(2));
+ results.push_back(inlinePtx.getResult(3));
+
+ mlir::Type vecF32Ty = mlir::VectorType::get({extent}, f32Ty);
+ mlir::Value undef = mlir::LLVM::UndefOp::create(builder, loc, vecF32Ty);
+ mlir::Type i32Ty = builder.getI32Type();
+ for (unsigned i = 0; i < extent; ++i)
+ undef = mlir::LLVM::InsertElementOp::create(
+ builder, loc, undef, results[i],
+ builder.createIntegerConstant(loc, i32Ty, i));
+
+ auto i128Ty = builder.getIntegerType(128);
+ auto i128VecTy = mlir::VectorType::get({1}, i128Ty);
+ mlir::Value vec128 =
+ mlir::vector::BitCastOp::create(builder, loc, i128VecTy, undef);
+ return mlir::vector::ExtractOp::create(builder, loc, vec128,
+ mlir::ArrayRef<int64_t>{0});
+}
+
mlir::Value
CUDAIntrinsicLibrary::genAtomicAnd(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
@@ -891,9 +959,8 @@ CUDAIntrinsicLibrary::genBarrierArrive(mlir::Type resultType,
assert(args.size() == 1);
mlir::Value barrier = convertPtrToNVVMSpace(
builder, loc, args[0], mlir::NVVM::NVVMMemorySpace::Shared);
- return mlir::NVVM::MBarrierArriveSharedOp::create(builder, loc, resultType,
- barrier)
- .getResult();
+ return mlir::NVVM::MBarrierArriveOp::create(builder, loc, resultType, barrier)
+ .getResult(0);
}
// BARRIER_ARRIBVE_CNT
@@ -982,6 +1049,60 @@ CUDAIntrinsicLibrary::genBarrierTryWaitSleep(mlir::Type resultType,
.getResult(0);
}
+static void insertValueAtPos(fir::FirOpBuilder &builder, mlir::Location loc,
+ fir::RecordType recTy, mlir::Value base,
+ mlir::Value dim, unsigned fieldPos) {
+ auto fieldName = recTy.getTypeList()[fieldPos].first;
+ mlir::Type fieldTy = recTy.getTypeList()[fieldPos].second;
+ mlir::Type fieldIndexType = fir::FieldType::get(base.getContext());
+ mlir::Value fieldIndex =
+ fir::FieldIndexOp::create(builder, loc, fieldIndexType, fieldName, recTy,
+ /*typeParams=*/mlir::ValueRange{});
+ mlir::Value coord = fir::CoordinateOp::create(
+ builder, loc, builder.getRefType(fieldTy), base, fieldIndex);
+ fir::StoreOp::create(builder, loc, dim, coord);
+}
+
+// CLUSTER_BLOCK_INDEX
+mlir::Value
+CUDAIntrinsicLibrary::genClusterBlockIndex(mlir::Type resultType,
+ llvm::ArrayRef<mlir::Value> args) {
+ assert(args.size() == 0);
+ auto recTy = mlir::cast<fir::RecordType>(resultType);
+ assert(recTy && "RecordType expepected");
+ mlir::Value res = fir::AllocaOp::create(builder, loc, resultType);
+ mlir::Type i32Ty = builder.getI32Type();
+ mlir::Value x = mlir::NVVM::BlockInClusterIdXOp::create(builder, loc, i32Ty);
+ mlir::Value one = builder.createIntegerConstant(loc, i32Ty, 1);
+ x = mlir::arith::AddIOp::create(builder, loc, x, one);
+ insertValueAtPos(builder, loc, recTy, res, x, 0);
+ mlir::Value y = mlir::NVVM::BlockInClusterIdYOp::create(builder, loc, i32Ty);
+ y = mlir::arith::AddIOp::create(builder, loc, y, one);
+ insertValueAtPos(builder, loc, recTy, res, y, 1);
+ mlir::Value z = mlir::NVVM::BlockInClusterIdZOp::create(builder, loc, i32Ty);
+ z = mlir::arith::AddIOp::create(builder, loc, z, one);
+ insertValueAtPos(builder, loc, recTy, res, z, 2);
+ return res;
+}
+
+// CLUSTER_DIM_BLOCKS
+mlir::Value
+CUDAIntrinsicLibrary::genClusterDimBlocks(mlir::Type resultType,
+ llvm::ArrayRef<mlir::Value> args) {
+ assert(args.size() == 0);
+ auto recTy = mlir::cast<fir::RecordType>(resultType);
+ assert(recTy && "RecordType expepected");
+ mlir::Value res = fir::AllocaOp::create(builder, loc, resultType);
+ mlir::Type i32Ty = builder.getI32Type();
+ mlir::Value x = mlir::NVVM::ClusterDimBlocksXOp::create(builder, loc, i32Ty);
+ insertValueAtPos(builder, loc, recTy, res, x, 0);
+ mlir::Value y = mlir::NVVM::ClusterDimBlocksYOp::create(builder, loc, i32Ty);
+ insertValueAtPos(builder, loc, recTy, res, y, 1);
+ mlir::Value z = mlir::NVVM::ClusterDimBlocksZOp::create(builder, loc, i32Ty);
+ insertValueAtPos(builder, loc, recTy, res, z, 2);
+ return res;
+}
+
// FENCE_PROXY_ASYNC
void CUDAIntrinsicLibrary::genFenceProxyAsync(
llvm::ArrayRef<fir::ExtendedValue> args) {
@@ -1081,55 +1202,84 @@ void CUDAIntrinsicLibrary::genSyncThreads(
mlir::Value
CUDAIntrinsicLibrary::genSyncThreadsAnd(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
- constexpr llvm::StringLiteral funcName = "llvm.nvvm.barrier0.and";
- mlir::MLIRContext *context = builder.getContext();
- mlir::Type i32 = builder.getI32Type();
- mlir::FunctionType ftype =
- mlir::FunctionType::get(context, {resultType}, {i32});
- auto funcOp = builder.createFunction(loc, funcName, ftype);
- mlir::Value arg = builder.createConvert(loc, i32, args[0]);
- return fir::CallOp::create(builder, loc, funcOp, {arg}).getResult(0);
+ mlir::Value arg = builder.createConvert(loc, builder.getI32Type(), args[0]);
+ return mlir::NVVM::BarrierOp::create(
+ builder, loc, resultType, {}, {},
+ mlir::NVVM::BarrierReductionAttr::get(
+ builder.getContext(), mlir::NVVM::BarrierReduction::AND),
+ arg)
+ .getResult(0);
}
// SYNCTHREADS_COUNT
mlir::Value
CUDAIntrinsicLibrary::genSyncThreadsCount(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
- constexpr llvm::StringLiteral funcName = "llvm.nvvm.barrier0.popc";
- mlir::MLIRContext *context = builder.getContext();
- mlir::Type i32 = builder.getI32Type();
- mlir::FunctionType ftype =
- mlir::FunctionType::get(context, {resultType}, {i32});
- auto funcOp = builder.createFunction(loc, funcName, ftype);
- mlir::Value arg = builder.createConvert(loc, i32, args[0]);
- return fir::CallOp::create(builder, loc, funcOp, {arg}).getResult(0);
+ mlir::Value arg = builder.createConvert(loc, builder.getI32Type(), args[0]);
+ return mlir::NVVM::BarrierOp::create(
+ builder, loc, resultType, {}, {},
+ mlir::NVVM::BarrierReductionAttr::get(
+ builder.getContext(), mlir::NVVM::BarrierReduction::POPC),
+ arg)
+ .getResult(0);
}
// SYNCTHREADS_OR
mlir::Value
CUDAIntrinsicLibrary::genSyncThreadsOr(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
- constexpr llvm::StringLiteral funcName = "llvm.nvvm.barrier0.or";
- mlir::MLIRContext *context = builder.getContext();
- mlir::Type i32 = builder.getI32Type();
- mlir::FunctionType ftype =
- mlir::FunctionType::get(context, {resultType}, {i32});
- auto funcOp = builder.createFunction(loc, funcName, ftype);
- mlir::Value arg = builder.createConvert(loc, i32, args[0]);
- return fir::CallOp::create(builder, loc, funcOp, {arg}).getResult(0);
+ mlir::Value arg = builder.createConvert(loc, builder.getI32Type(), args[0]);
+ return mlir::NVVM::BarrierOp::create(
+ builder, loc, resultType, {}, {},
+ mlir::NVVM::BarrierReductionAttr::get(
+ builder.getContext(), mlir::NVVM::BarrierReduction::OR),
+ arg)
+ .getResult(0);
}
// SYNCWARP
void CUDAIntrinsicLibrary::genSyncWarp(
llvm::ArrayRef<fir::ExtendedValue> args) {
assert(args.size() == 1);
- constexpr llvm::StringLiteral funcName = "llvm.nvvm.bar.warp.sync";
- mlir::Value mask = fir::getBase(args[0]);
- mlir::FunctionType funcType =
- mlir::FunctionType::get(builder.getContext(), {mask.getType()}, {});
- auto funcOp = builder.createFunction(loc, funcName, funcType);
- llvm::SmallVector<mlir::Value> argsList{mask};
- fir::CallOp::create(builder, loc, funcOp, argsList);
+ mlir::NVVM::SyncWarpOp::create(builder, loc, fir::getBase(args[0]));
+}
+
+// THIS_CLUSTER
+mlir::Value
+CUDAIntrinsicLibrary::genThisCluster(mlir::Type resultType,
+ llvm::ArrayRef<mlir::Value> args) {
+ assert(args.size() == 0);
+ auto recTy = mlir::cast<fir::RecordType>(resultType);
+ assert(recTy && "RecordType expepected");
+ mlir::Value res = fir::AllocaOp::create(builder, loc, resultType);
+ mlir::Type i32Ty = builder.getI32Type();
+
+ // SIZE
+ mlir::Value size = mlir::NVVM::ClusterDim::create(builder, loc, i32Ty);
+ auto sizeFieldName = recTy.getTypeList()[1].first;
+ mlir::Type sizeFieldTy = recTy.getTypeList()[1].second;
+ mlir::Type fieldIndexType = fir::FieldType::get(resultType.getContext());
+ mlir::Value sizeFieldIndex = fir::FieldIndexOp::create(
+ builder, loc, fieldIndexType, sizeFieldName, recTy,
+ /*typeParams=*/mlir::ValueRange{});
+ mlir::Value sizeCoord = fir::CoordinateOp::create(
+ builder, loc, builder.getRefType(sizeFieldTy), res, sizeFieldIndex);
+ fir::StoreOp::create(builder, loc, size, sizeCoord);
+
+ // RANK
+ mlir::Value rank = mlir::NVVM::ClusterId::create(builder, loc, i32Ty);
+ mlir::Value one = builder.createIntegerConstant(loc, i32Ty, 1);
+ rank = mlir::arith::AddIOp::create(builder, loc, rank, one);
+ auto rankFieldName = recTy.getTypeList()[2].first;
+ mlir::Type rankFieldTy = recTy.getTypeList()[2].second;
+ mlir::Value rankFieldIndex = fir::FieldIndexOp::create(
+ builder, loc, fieldIndexType, rankFieldName, recTy,
+ /*typeParams=*/mlir::ValueRange{});
+ mlir::Value rankCoord = fir::CoordinateOp::create(
+ builder, loc, builder.getRefType(rankFieldTy), res, rankFieldIndex);
+ fir::StoreOp::create(builder, loc, rank, rankCoord);
+
+ return res;
}
// THIS_GRID
@@ -1313,37 +1463,12 @@ CUDAIntrinsicLibrary::genThisWarp(mlir::Type resultType,
return res;
}
-// THREADFENCE
+// THREADFENCE, THREADFENCE_BLOCK, THREADFENCE_SYSTEM
+template <mlir::NVVM::MemScopeKind scope>
void CUDAIntrinsicLibrary::genThreadFence(
llvm::ArrayRef<fir::ExtendedValue> args) {
- constexpr llvm::StringLiteral funcName = "llvm.nvvm.membar.gl";
- mlir::FunctionType funcType =
- mlir::FunctionType::get(builder.getContext(), {}, {});
- auto funcOp = builder.createFunction(loc, funcName, funcType);
- llvm::SmallVector<mlir::Value> noArgs;
- fir::CallOp::create(builder, loc, funcOp, noArgs);
-}
-
-// THREADFENCE_BLOCK
-void CUDAIntrinsicLibrary::genThreadFenceBlock(
- llvm::ArrayRef<fir::ExtendedValue> args) {
- constexpr llvm::StringLiteral funcName = "llvm.nvvm.membar.cta";
- mlir::FunctionType funcType =
- mlir::FunctionType::get(builder.getContext(), {}, {});
- auto funcOp = builder.createFunction(loc, funcName, funcType);
- llvm::SmallVector<mlir::Value> noArgs;
- fir::CallOp::create(builder, loc, funcOp, noArgs);
-}
-
-// THREADFENCE_SYSTEM
-void CUDAIntrinsicLibrary::genThreadFenceSystem(
- llvm::ArrayRef<fir::ExtendedValue> args) {
- constexpr llvm::StringLiteral funcName = "llvm.nvvm.membar.sys";
- mlir::FunctionType funcType =
- mlir::FunctionType::get(builder.getContext(), {}, {});
- auto funcOp = builder.createFunction(loc, funcName, funcType);
- llvm::SmallVector<mlir::Value> noArgs;
- fir::CallOp::create(builder, loc, funcOp, noArgs);
+ assert(args.size() == 0);
+ mlir::NVVM::MembarOp::create(builder, loc, scope);
}
// TMA_BULK_COMMIT_GROUP
@@ -1368,6 +1493,13 @@ void CUDAIntrinsicLibrary::genTMABulkG2S(
builder, loc, dst, src, barrier, fir::getBase(args[3]), {}, {});
}
+static void setAlignment(mlir::Value ptr, unsigned alignment) {
+ if (auto declareOp = mlir::dyn_cast<hlfir::DeclareOp>(ptr.getDefiningOp()))
+ if (auto sharedOp = mlir::dyn_cast<cuf::SharedMemoryOp>(
+ declareOp.getMemref().getDefiningOp()))
+ sharedOp.setAlignment(alignment);
+}
+
static void genTMABulkLoad(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Value barrier, mlir::Value src,
mlir::Value dst, mlir::Value nelem,
@@ -1375,6 +1507,7 @@ static void genTMABulkLoad(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Value size = mlir::arith::MulIOp::create(builder, loc, nelem, eleSize);
auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(builder.getContext());
barrier = builder.createConvert(loc, llvmPtrTy, barrier);
+ setAlignment(dst, kTMAAlignment);
dst = builder.createConvert(loc, llvmPtrTy, dst);
src = builder.createConvert(loc, llvmPtrTy, src);
mlir::NVVM::InlinePtxOp::create(
@@ -1478,6 +1611,7 @@ static void genTMABulkStore(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Value src, mlir::Value dst, mlir::Value count,
mlir::Value eleSize) {
mlir::Value size = mlir::arith::MulIOp::create(builder, loc, eleSize, count);
+ setAlignment(src, kTMAAlignment);
src = convertPtrToNVVMSpace(builder, loc, src,
mlir::NVVM::NVVMMemorySpace::Shared);
dst = convertPtrToNVVMSpace(builder, loc, dst,
diff --git a/flang/lib/Optimizer/Builder/CUFCommon.cpp b/flang/lib/Optimizer/Builder/CUFCommon.cpp
index cf7588f..2266f4d 100644
--- a/flang/lib/Optimizer/Builder/CUFCommon.cpp
+++ b/flang/lib/Optimizer/Builder/CUFCommon.cpp
@@ -9,6 +9,7 @@
#include "flang/Optimizer/Builder/CUFCommon.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
+#include "flang/Optimizer/Dialect/Support/KindMapping.h"
#include "flang/Optimizer/HLFIR/HLFIROps.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
@@ -91,3 +92,66 @@ void cuf::genPointerSync(const mlir::Value box, fir::FirOpBuilder &builder) {
}
}
}
+
+int cuf::computeElementByteSize(mlir::Location loc, mlir::Type type,
+ fir::KindMapping &kindMap,
+ bool emitErrorOnFailure) {
+ auto eleTy = fir::unwrapSequenceType(type);
+ if (auto t{mlir::dyn_cast<mlir::IntegerType>(eleTy)})
+ return t.getWidth() / 8;
+ if (auto t{mlir::dyn_cast<mlir::FloatType>(eleTy)})
+ return t.getWidth() / 8;
+ if (auto t{mlir::dyn_cast<fir::LogicalType>(eleTy)})
+ return kindMap.getLogicalBitsize(t.getFKind()) / 8;
+ if (auto t{mlir::dyn_cast<mlir::ComplexType>(eleTy)}) {
+ int elemSize =
+ mlir::cast<mlir::FloatType>(t.getElementType()).getWidth() / 8;
+ return 2 * elemSize;
+ }
+ if (auto t{mlir::dyn_cast<fir::CharacterType>(eleTy)})
+ return kindMap.getCharacterBitsize(t.getFKind()) / 8;
+ if (emitErrorOnFailure)
+ mlir::emitError(loc, "unsupported type");
+ return 0;
+}
+
+mlir::Value cuf::computeElementCount(mlir::PatternRewriter &rewriter,
+ mlir::Location loc,
+ mlir::Value shapeOperand,
+ mlir::Type seqType,
+ mlir::Type targetType) {
+ if (shapeOperand) {
+ // Dynamic extent - extract from shape operand
+ llvm::SmallVector<mlir::Value> extents;
+ if (auto shapeOp =
+ mlir::dyn_cast<fir::ShapeOp>(shapeOperand.getDefiningOp())) {
+ extents = shapeOp.getExtents();
+ } else if (auto shapeShiftOp = mlir::dyn_cast<fir::ShapeShiftOp>(
+ shapeOperand.getDefiningOp())) {
+ for (auto i : llvm::enumerate(shapeShiftOp.getPairs()))
+ if (i.index() & 1)
+ extents.push_back(i.value());
+ }
+
+ if (extents.empty())
+ return mlir::Value();
+
+ // Compute total element count by multiplying all dimensions
+ mlir::Value count =
+ fir::ConvertOp::create(rewriter, loc, targetType, extents[0]);
+ for (unsigned i = 1; i < extents.size(); ++i) {
+ auto operand =
+ fir::ConvertOp::create(rewriter, loc, targetType, extents[i]);
+ count = mlir::arith::MulIOp::create(rewriter, loc, count, operand);
+ }
+ return count;
+ } else {
+ // Static extent - use constant array size
+ if (auto seqTy = mlir::dyn_cast_or_null<fir::SequenceType>(seqType)) {
+ mlir::IntegerAttr attr =
+ rewriter.getIntegerAttr(targetType, seqTy.getConstantArraySize());
+ return mlir::arith::ConstantOp::create(rewriter, loc, targetType, attr);
+ }
+ }
+ return mlir::Value();
+}
diff --git a/flang/lib/Optimizer/Builder/FIRBuilder.cpp b/flang/lib/Optimizer/Builder/FIRBuilder.cpp
index 5da27d1..c704ac7 100644
--- a/flang/lib/Optimizer/Builder/FIRBuilder.cpp
+++ b/flang/lib/Optimizer/Builder/FIRBuilder.cpp
@@ -427,7 +427,8 @@ mlir::Value fir::FirOpBuilder::genTempDeclareOp(
builder, loc, memref.getType(), memref, shape, typeParams,
/*dummy_scope=*/nullptr,
/*storage=*/nullptr,
- /*storage_offset=*/0, nameAttr, fortranAttrs, cuf::DataAttributeAttr{});
+ /*storage_offset=*/0, nameAttr, fortranAttrs, cuf::DataAttributeAttr{},
+ /*dummy_arg_no=*/mlir::IntegerAttr{});
}
mlir::Value fir::FirOpBuilder::genStackSave(mlir::Location loc) {
@@ -1392,12 +1393,10 @@ fir::ExtendedValue fir::factory::arraySectionElementToExtendedValue(
return fir::factory::componentToExtendedValue(builder, loc, element);
}
-void fir::factory::genScalarAssignment(fir::FirOpBuilder &builder,
- mlir::Location loc,
- const fir::ExtendedValue &lhs,
- const fir::ExtendedValue &rhs,
- bool needFinalization,
- bool isTemporaryLHS) {
+void fir::factory::genScalarAssignment(
+ fir::FirOpBuilder &builder, mlir::Location loc,
+ const fir::ExtendedValue &lhs, const fir::ExtendedValue &rhs,
+ bool needFinalization, bool isTemporaryLHS, mlir::ArrayAttr accessGroups) {
assert(lhs.rank() == 0 && rhs.rank() == 0 && "must be scalars");
auto type = fir::unwrapSequenceType(
fir::unwrapPassByRefType(fir::getBase(lhs).getType()));
@@ -1419,7 +1418,9 @@ void fir::factory::genScalarAssignment(fir::FirOpBuilder &builder,
mlir::Value lhsAddr = fir::getBase(lhs);
rhsVal = builder.createConvert(loc, fir::unwrapRefType(lhsAddr.getType()),
rhsVal);
- fir::StoreOp::create(builder, loc, rhsVal, lhsAddr);
+ fir::StoreOp store = fir::StoreOp::create(builder, loc, rhsVal, lhsAddr);
+ if (accessGroups)
+ store.setAccessGroupsAttr(accessGroups);
}
}
@@ -1670,6 +1671,26 @@ mlir::Value fir::factory::createZeroValue(fir::FirOpBuilder &builder,
"numeric or logical type");
}
+mlir::Value fir::factory::createOneValue(fir::FirOpBuilder &builder,
+ mlir::Location loc, mlir::Type type) {
+ mlir::Type i1 = builder.getIntegerType(1);
+ if (mlir::isa<fir::LogicalType>(type) || type == i1)
+ return builder.createConvert(loc, type, builder.createBool(loc, true));
+ if (fir::isa_integer(type))
+ return builder.createIntegerConstant(loc, type, 1);
+ if (fir::isa_real(type))
+ return builder.createRealOneConstant(loc, type);
+ if (fir::isa_complex(type)) {
+ fir::factory::Complex complexHelper(builder, loc);
+ mlir::Type partType = complexHelper.getComplexPartType(type);
+ mlir::Value realPart = builder.createRealOneConstant(loc, partType);
+ mlir::Value imagPart = builder.createRealZeroConstant(loc, partType);
+ return complexHelper.createComplex(type, realPart, imagPart);
+ }
+ fir::emitFatalError(loc, "internal: trying to generate one value of non "
+ "numeric or logical type");
+}
+
std::optional<std::int64_t>
fir::factory::getExtentFromTriplet(mlir::Value lb, mlir::Value ub,
mlir::Value stride) {
diff --git a/flang/lib/Optimizer/Builder/HLFIRTools.cpp b/flang/lib/Optimizer/Builder/HLFIRTools.cpp
index 93dfc57..a345dcb 100644
--- a/flang/lib/Optimizer/Builder/HLFIRTools.cpp
+++ b/flang/lib/Optimizer/Builder/HLFIRTools.cpp
@@ -250,7 +250,7 @@ hlfir::genDeclare(mlir::Location loc, fir::FirOpBuilder &builder,
const fir::ExtendedValue &exv, llvm::StringRef name,
fir::FortranVariableFlagsAttr flags, mlir::Value dummyScope,
mlir::Value storage, std::uint64_t storageOffset,
- cuf::DataAttributeAttr dataAttr) {
+ cuf::DataAttributeAttr dataAttr, unsigned dummyArgNo) {
mlir::Value base = fir::getBase(exv);
assert(fir::conformsWithPassByRef(base.getType()) &&
@@ -281,7 +281,7 @@ hlfir::genDeclare(mlir::Location loc, fir::FirOpBuilder &builder,
[](const auto &) {});
auto declareOp = hlfir::DeclareOp::create(
builder, loc, base, name, shapeOrShift, lenParams, dummyScope, storage,
- storageOffset, flags, dataAttr);
+ storageOffset, flags, dataAttr, dummyArgNo);
return mlir::cast<fir::FortranVariableOpInterface>(declareOp.getOperation());
}
@@ -402,9 +402,9 @@ hlfir::Entity hlfir::genVariableBox(mlir::Location loc,
fir::BoxType::get(var.getElementOrSequenceType(), isVolatile);
if (forceBoxType) {
boxType = forceBoxType;
- mlir::Type baseType =
- fir::ReferenceType::get(fir::unwrapRefType(forceBoxType.getEleTy()));
- addr = builder.createConvert(loc, baseType, addr);
+ mlir::Type baseType = fir::ReferenceType::get(
+ fir::unwrapRefType(forceBoxType.getEleTy()), forceBoxType.isVolatile());
+ addr = builder.createConvertWithVolatileCast(loc, baseType, addr);
}
auto embox = fir::EmboxOp::create(builder, loc, boxType, addr, shape,
/*slice=*/mlir::Value{}, typeParams);
@@ -1392,6 +1392,66 @@ bool hlfir::elementalOpMustProduceTemp(hlfir::ElementalOp elemental) {
return false;
}
+static void combineAndStoreElement(
+ mlir::Location loc, fir::FirOpBuilder &builder, hlfir::Entity lhs,
+ hlfir::Entity rhs, bool temporaryLHS,
+ std::function<hlfir::Entity(mlir::Location, fir::FirOpBuilder &,
+ hlfir::Entity, hlfir::Entity)> *combiner) {
+ hlfir::Entity valueToAssign = hlfir::loadTrivialScalar(loc, builder, rhs);
+ if (combiner) {
+ hlfir::Entity lhsValue = hlfir::loadTrivialScalar(loc, builder, lhs);
+ valueToAssign = (*combiner)(loc, builder, lhsValue, valueToAssign);
+ }
+ hlfir::AssignOp::create(builder, loc, valueToAssign, lhs,
+ /*realloc=*/false,
+ /*keep_lhs_length_if_realloc=*/false,
+ /*temporary_lhs=*/temporaryLHS);
+}
+
+void hlfir::genNoAliasArrayAssignment(
+ mlir::Location loc, fir::FirOpBuilder &builder, hlfir::Entity rhs,
+ hlfir::Entity lhs, bool emitWorkshareLoop, bool temporaryLHS,
+ std::function<hlfir::Entity(mlir::Location, fir::FirOpBuilder &,
+ hlfir::Entity, hlfir::Entity)> *combiner) {
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ rhs = hlfir::derefPointersAndAllocatables(loc, builder, rhs);
+ lhs = hlfir::derefPointersAndAllocatables(loc, builder, lhs);
+ mlir::Value lhsShape = hlfir::genShape(loc, builder, lhs);
+ llvm::SmallVector<mlir::Value> lhsExtents =
+ hlfir::getIndexExtents(loc, builder, lhsShape);
+ mlir::Value rhsShape = hlfir::genShape(loc, builder, rhs);
+ llvm::SmallVector<mlir::Value> rhsExtents =
+ hlfir::getIndexExtents(loc, builder, rhsShape);
+ llvm::SmallVector<mlir::Value> extents =
+ fir::factory::deduceOptimalExtents(lhsExtents, rhsExtents);
+ hlfir::LoopNest loopNest =
+ hlfir::genLoopNest(loc, builder, extents,
+ /*isUnordered=*/true, emitWorkshareLoop);
+ builder.setInsertionPointToStart(loopNest.body);
+ auto rhsArrayElement =
+ hlfir::getElementAt(loc, builder, rhs, loopNest.oneBasedIndices);
+ rhsArrayElement = hlfir::loadTrivialScalar(loc, builder, rhsArrayElement);
+ auto lhsArrayElement =
+ hlfir::getElementAt(loc, builder, lhs, loopNest.oneBasedIndices);
+ combineAndStoreElement(loc, builder, lhsArrayElement, rhsArrayElement,
+ temporaryLHS, combiner);
+}
+
+void hlfir::genNoAliasAssignment(
+ mlir::Location loc, fir::FirOpBuilder &builder, hlfir::Entity rhs,
+ hlfir::Entity lhs, bool emitWorkshareLoop, bool temporaryLHS,
+ std::function<hlfir::Entity(mlir::Location, fir::FirOpBuilder &,
+ hlfir::Entity, hlfir::Entity)> *combiner) {
+ if (lhs.isArray()) {
+ genNoAliasArrayAssignment(loc, builder, rhs, lhs, emitWorkshareLoop,
+ temporaryLHS, combiner);
+ return;
+ }
+ rhs = hlfir::derefPointersAndAllocatables(loc, builder, rhs);
+ lhs = hlfir::derefPointersAndAllocatables(loc, builder, lhs);
+ combineAndStoreElement(loc, builder, lhs, rhs, temporaryLHS, combiner);
+}
+
std::pair<hlfir::Entity, bool>
hlfir::createTempFromMold(mlir::Location loc, fir::FirOpBuilder &builder,
hlfir::Entity mold) {
diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index 3eb6044..75a74ee 100644
--- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
@@ -91,6 +91,11 @@ static bool isStaticallyAbsent(llvm::ArrayRef<mlir::Value> args,
size_t argIndex) {
return args.size() <= argIndex || !args[argIndex];
}
+static bool isOptional(mlir::Value value) {
+ auto varIface = mlir::dyn_cast_or_null<fir::FortranVariableOpInterface>(
+ value.getDefiningOp());
+ return varIface && varIface.isOptional();
+}
/// Test if an ExtendedValue is present. This is used to test if an intrinsic
/// argument is present at compile time. This does not imply that the related
@@ -303,6 +308,10 @@ static constexpr IntrinsicHandler handlers[]{
{"back", asValue, handleDynamicOptional}}},
/*isElemental=*/false},
{"floor", &I::genFloor},
+ {"flush",
+ &I::genFlush,
+ {{{"unit", asAddr}}},
+ /*isElemental=*/false},
{"fraction", &I::genFraction},
{"free", &I::genFree},
{"fseek",
@@ -340,6 +349,10 @@ static constexpr IntrinsicHandler handlers[]{
{"trim_name", asAddr, handleDynamicOptional},
{"errmsg", asBox, handleDynamicOptional}}},
/*isElemental=*/false},
+ {"get_team",
+ &I::genGetTeam,
+ {{{"level", asValue, handleDynamicOptional}}},
+ /*isElemental=*/false},
{"getcwd",
&I::genGetCwd,
{{{"c", asBox}, {"status", asAddr, handleDynamicOptional}}},
@@ -486,6 +499,10 @@ static constexpr IntrinsicHandler handlers[]{
{"dim", asValue},
{"mask", asBox, handleDynamicOptional}}},
/*isElemental=*/false},
+ {"irand",
+ &I::genIrand,
+ {{{"i", asAddr, handleDynamicOptional}}},
+ /*isElemental=*/false},
{"is_contiguous",
&I::genIsContiguous,
{{{"array", asBox}}},
@@ -612,6 +629,10 @@ static constexpr IntrinsicHandler handlers[]{
&I::genPutenv,
{{{"str", asAddr}, {"status", asAddr, handleDynamicOptional}}},
/*isElemental=*/false},
+ {"rand",
+ &I::genRand,
+ {{{"i", asAddr, handleDynamicOptional}}},
+ /*isElemental=*/false},
{"random_init",
&I::genRandomInit,
{{{"repeatable", asValue}, {"image_distinct", asValue}}},
@@ -706,6 +727,10 @@ static constexpr IntrinsicHandler handlers[]{
{"shifta", &I::genShiftA},
{"shiftl", &I::genShift<mlir::arith::ShLIOp>},
{"shiftr", &I::genShift<mlir::arith::ShRUIOp>},
+ {"show_descriptor",
+ &I::genShowDescriptor,
+ {{{"d", asBox}}},
+ /*isElemental=*/false},
{"sign", &I::genSign},
{"signal",
&I::genSignalSubroutine,
@@ -749,6 +774,10 @@ static constexpr IntrinsicHandler handlers[]{
/*isElemental=*/false},
{"tand", &I::genTand},
{"tanpi", &I::genTanpi},
+ {"team_number",
+ &I::genTeamNumber,
+ {{{"team", asBox, handleDynamicOptional}}},
+ /*isElemental=*/false},
{"this_image",
&I::genThisImage,
{{{"coarray", asBox},
@@ -3934,6 +3963,40 @@ mlir::Value IntrinsicLibrary::genFloor(mlir::Type resultType,
return builder.createConvert(loc, resultType, floor);
}
+// FLUSH
+void IntrinsicLibrary::genFlush(llvm::ArrayRef<fir::ExtendedValue> args) {
+ assert(args.size() == 1);
+
+ mlir::Value unit;
+ if (isStaticallyAbsent(args[0]))
+ // Give a sentinal value of `-1` on the `()` case.
+ unit = builder.createIntegerConstant(loc, builder.getI32Type(), -1);
+ else {
+ unit = fir::getBase(args[0]);
+ if (isOptional(unit)) {
+ mlir::Value isPresent =
+ fir::IsPresentOp::create(builder, loc, builder.getI1Type(), unit);
+ unit = builder
+ .genIfOp(loc, builder.getI32Type(), isPresent,
+ /*withElseRegion=*/true)
+ .genThen([&]() {
+ mlir::Value loaded = fir::LoadOp::create(builder, loc, unit);
+ fir::ResultOp::create(builder, loc, loaded);
+ })
+ .genElse([&]() {
+ mlir::Value negOne = builder.createIntegerConstant(
+ loc, builder.getI32Type(), -1);
+ fir::ResultOp::create(builder, loc, negOne);
+ })
+ .getResults()[0];
+ } else {
+ unit = fir::LoadOp::create(builder, loc, unit);
+ }
+ }
+
+ fir::runtime::genFlush(builder, loc, unit);
+}
+
// FRACTION
mlir::Value IntrinsicLibrary::genFraction(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
@@ -4013,6 +4076,15 @@ IntrinsicLibrary::genFtell(std::optional<mlir::Type> resultType,
}
}
+// GET_TEAM
+mlir::Value IntrinsicLibrary::genGetTeam(mlir::Type resultType,
+ llvm::ArrayRef<mlir::Value> args) {
+ converter->checkCoarrayEnabled();
+ assert(args.size() == 1);
+ return mif::GetTeamOp::create(builder, loc, fir::BoxType::get(resultType),
+ /*level*/ args[0]);
+}
+
// GETCWD
fir::ExtendedValue
IntrinsicLibrary::genGetCwd(std::optional<mlir::Type> resultType,
@@ -6098,6 +6170,20 @@ IntrinsicLibrary::genIparity(mlir::Type resultType,
"IPARITY", resultType, args);
}
+// IRAND
+fir::ExtendedValue
+IntrinsicLibrary::genIrand(mlir::Type resultType,
+ llvm::ArrayRef<fir::ExtendedValue> args) {
+ assert(args.size() == 1);
+ mlir::Value i =
+ isStaticallyPresent(args[0])
+ ? fir::getBase(args[0])
+ : fir::AbsentOp::create(builder, loc,
+ builder.getRefType(builder.getI32Type()))
+ .getResult();
+ return fir::runtime::genIrand(builder, loc, i);
+}
+
// IS_CONTIGUOUS
fir::ExtendedValue
IntrinsicLibrary::genIsContiguous(mlir::Type resultType,
@@ -6281,12 +6367,6 @@ IntrinsicLibrary::genCharacterCompare(mlir::Type resultType,
fir::getBase(args[1]), fir::getLen(args[1]));
}
-static bool isOptional(mlir::Value value) {
- auto varIface = mlir::dyn_cast_or_null<fir::FortranVariableOpInterface>(
- value.getDefiningOp());
- return varIface && varIface.isOptional();
-}
-
// LOC
fir::ExtendedValue
IntrinsicLibrary::genLoc(mlir::Type resultType,
@@ -6509,11 +6589,9 @@ static mlir::Value genFastMod(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Value IntrinsicLibrary::genMod(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
auto mod = builder.getModule();
- bool dontUseFastRealMod = false;
- bool canUseApprox = mlir::arith::bitEnumContainsAny(
- builder.getFastMathFlags(), mlir::arith::FastMathFlags::afn);
- if (auto attr = mod->getAttrOfType<mlir::BoolAttr>("fir.no_fast_real_mod"))
- dontUseFastRealMod = attr.getValue();
+ bool useFastRealMod = false;
+ if (auto attr = mod->getAttrOfType<mlir::BoolAttr>("fir.fast_real_mod"))
+ useFastRealMod = attr.getValue();
assert(args.size() == 2);
if (resultType.isUnsignedInteger()) {
@@ -6526,7 +6604,7 @@ mlir::Value IntrinsicLibrary::genMod(mlir::Type resultType,
if (mlir::isa<mlir::IntegerType>(resultType))
return mlir::arith::RemSIOp::create(builder, loc, args[0], args[1]);
- if (resultType.isFloat() && canUseApprox && !dontUseFastRealMod) {
+ if (resultType.isFloat() && useFastRealMod) {
// Treat MOD as an approximate function and code-gen inline code
// instead of calling into the Fortran runtime library.
return builder.createConvert(loc, resultType,
@@ -7132,6 +7210,19 @@ IntrinsicLibrary::genPutenv(std::optional<mlir::Type> resultType,
return {};
}
+// RAND
+fir::ExtendedValue
+IntrinsicLibrary::genRand(mlir::Type, llvm::ArrayRef<fir::ExtendedValue> args) {
+ assert(args.size() == 1);
+ mlir::Value i =
+ isStaticallyPresent(args[0])
+ ? fir::getBase(args[0])
+ : fir::AbsentOp::create(builder, loc,
+ builder.getRefType(builder.getI32Type()))
+ .getResult();
+ return fir::runtime::genRand(builder, loc, i);
+}
+
// RANDOM_INIT
void IntrinsicLibrary::genRandomInit(llvm::ArrayRef<fir::ExtendedValue> args) {
assert(args.size() == 2);
@@ -7797,6 +7888,16 @@ mlir::Value IntrinsicLibrary::genShiftA(mlir::Type resultType,
return result;
}
+void IntrinsicLibrary::genShowDescriptor(
+ llvm::ArrayRef<fir::ExtendedValue> args) {
+ assert(args.size() == 1 && "expected single argument for show_descriptor");
+ const mlir::Value descriptor = fir::getBase(args[0]);
+
+ assert(fir::isa_box_type(descriptor.getType()) &&
+ "argument must have been lowered to box type");
+ fir::runtime::genShowDescriptor(builder, loc, descriptor);
+}
+
// SIGNAL
void IntrinsicLibrary::genSignalSubroutine(
llvm::ArrayRef<fir::ExtendedValue> args) {
@@ -7953,6 +8054,16 @@ mlir::Value IntrinsicLibrary::genTanpi(mlir::Type resultType,
return getRuntimeCallGenerator("tan", ftype)(builder, loc, {arg});
}
+// TEAM_NUMBER
+fir::ExtendedValue
+IntrinsicLibrary::genTeamNumber(mlir::Type,
+ llvm::ArrayRef<fir::ExtendedValue> args) {
+ converter->checkCoarrayEnabled();
+ assert(args.size() == 1);
+ return mif::TeamNumberOp::create(builder, loc,
+ /*team*/ fir::getBase(args[0]));
+}
+
// THIS_IMAGE
fir::ExtendedValue
IntrinsicLibrary::genThisImage(mlir::Type resultType,
diff --git a/flang/lib/Optimizer/Builder/Runtime/Character.cpp b/flang/lib/Optimizer/Builder/Runtime/Character.cpp
index 540ecba..2f1772f 100644
--- a/flang/lib/Optimizer/Builder/Runtime/Character.cpp
+++ b/flang/lib/Optimizer/Builder/Runtime/Character.cpp
@@ -94,27 +94,34 @@ fir::runtime::genCharCompare(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::arith::CmpIPredicate cmp,
mlir::Value lhsBuff, mlir::Value lhsLen,
mlir::Value rhsBuff, mlir::Value rhsLen) {
- mlir::func::FuncOp beginFunc;
- switch (discoverKind(lhsBuff.getType())) {
+ int lhsKind = discoverKind(lhsBuff.getType());
+ int rhsKind = discoverKind(rhsBuff.getType());
+ if (lhsKind != rhsKind) {
+ fir::emitFatalError(loc, "runtime does not support comparison of different "
+ "CHARACTER kind values");
+ }
+ mlir::func::FuncOp func;
+ switch (lhsKind) {
case 1:
- beginFunc = fir::runtime::getRuntimeFunc<mkRTKey(CharacterCompareScalar1)>(
+ func = fir::runtime::getRuntimeFunc<mkRTKey(CharacterCompareScalar1)>(
loc, builder);
break;
case 2:
- beginFunc = fir::runtime::getRuntimeFunc<mkRTKey(CharacterCompareScalar2)>(
+ func = fir::runtime::getRuntimeFunc<mkRTKey(CharacterCompareScalar2)>(
loc, builder);
break;
case 4:
- beginFunc = fir::runtime::getRuntimeFunc<mkRTKey(CharacterCompareScalar4)>(
+ func = fir::runtime::getRuntimeFunc<mkRTKey(CharacterCompareScalar4)>(
loc, builder);
break;
default:
- llvm_unreachable("runtime does not support CHARACTER KIND");
+ fir::emitFatalError(
+ loc, "unsupported CHARACTER kind value. Runtime expects 1, 2, or 4.");
}
- auto fTy = beginFunc.getFunctionType();
+ auto fTy = func.getFunctionType();
auto args = fir::runtime::createArguments(builder, loc, fTy, lhsBuff, rhsBuff,
lhsLen, rhsLen);
- auto tri = fir::CallOp::create(builder, loc, beginFunc, args).getResult(0);
+ auto tri = fir::CallOp::create(builder, loc, func, args).getResult(0);
auto zero = builder.createIntegerConstant(loc, tri.getType(), 0);
return mlir::arith::CmpIOp::create(builder, loc, cmp, tri, zero);
}
diff --git a/flang/lib/Optimizer/Builder/Runtime/Intrinsics.cpp b/flang/lib/Optimizer/Builder/Runtime/Intrinsics.cpp
index 110b1b2..a5f16f8 100644
--- a/flang/lib/Optimizer/Builder/Runtime/Intrinsics.cpp
+++ b/flang/lib/Optimizer/Builder/Runtime/Intrinsics.cpp
@@ -137,6 +137,15 @@ void fir::runtime::genEtime(fir::FirOpBuilder &builder, mlir::Location loc,
fir::CallOp::create(builder, loc, runtimeFunc, args);
}
+void fir::runtime::genFlush(fir::FirOpBuilder &builder, mlir::Location loc,
+ mlir::Value unit) {
+ auto runtimeFunc = fir::runtime::getRuntimeFunc<mkRTKey(Flush)>(loc, builder);
+ llvm::SmallVector<mlir::Value> args = fir::runtime::createArguments(
+ builder, loc, runtimeFunc.getFunctionType(), unit);
+
+ fir::CallOp::create(builder, loc, runtimeFunc, args);
+}
+
void fir::runtime::genFree(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Value ptr) {
auto runtimeFunc = fir::runtime::getRuntimeFunc<mkRTKey(Free)>(loc, builder);
@@ -461,3 +470,34 @@ mlir::Value fir::runtime::genChdir(fir::FirOpBuilder &builder,
fir::runtime::createArguments(builder, loc, func.getFunctionType(), name);
return fir::CallOp::create(builder, loc, func, args).getResult(0);
}
+
+mlir::Value fir::runtime::genIrand(fir::FirOpBuilder &builder,
+ mlir::Location loc, mlir::Value i) {
+ auto runtimeFunc = fir::runtime::getRuntimeFunc<mkRTKey(Irand)>(loc, builder);
+ mlir::FunctionType runtimeFuncTy = runtimeFunc.getFunctionType();
+
+ llvm::SmallVector<mlir::Value> args =
+ fir::runtime::createArguments(builder, loc, runtimeFuncTy, i);
+ return fir::CallOp::create(builder, loc, runtimeFunc, args).getResult(0);
+}
+
+mlir::Value fir::runtime::genRand(fir::FirOpBuilder &builder,
+ mlir::Location loc, mlir::Value i) {
+ auto runtimeFunc = fir::runtime::getRuntimeFunc<mkRTKey(Rand)>(loc, builder);
+ mlir::FunctionType runtimeFuncTy = runtimeFunc.getFunctionType();
+
+ mlir::Value sourceFile = fir::factory::locationToFilename(builder, loc);
+ mlir::Value sourceLine =
+ fir::factory::locationToLineNo(builder, loc, runtimeFuncTy.getInput(2));
+
+ llvm::SmallVector<mlir::Value> args = fir::runtime::createArguments(
+ builder, loc, runtimeFuncTy, i, sourceFile, sourceLine);
+ return fir::CallOp::create(builder, loc, runtimeFunc, args).getResult(0);
+}
+
+void fir::runtime::genShowDescriptor(fir::FirOpBuilder &builder,
+ mlir::Location loc, mlir::Value descAddr) {
+ mlir::func::FuncOp func{
+ fir::runtime::getRuntimeFunc<mkRTKey(ShowDescriptor)>(loc, builder)};
+ fir::CallOp::create(builder, loc, func, descAddr);
+}
diff --git a/flang/lib/Optimizer/Builder/Runtime/Reduction.cpp b/flang/lib/Optimizer/Builder/Runtime/Reduction.cpp
index 157d435..343d848 100644
--- a/flang/lib/Optimizer/Builder/Runtime/Reduction.cpp
+++ b/flang/lib/Optimizer/Builder/Runtime/Reduction.cpp
@@ -1841,7 +1841,7 @@ mlir::Value fir::runtime::genReduce(fir::FirOpBuilder &builder,
assert((fir::isa_real(eleTy) || fir::isa_integer(eleTy) ||
mlir::isa<fir::LogicalType>(eleTy)) &&
- "expect real, interger or logical");
+ "expect real, integer or logical");
auto [cat, kind] = fir::mlirTypeToCategoryKind(loc, eleTy);
mlir::func::FuncOp func;
generated by cgit v1.1 at 2026-01-07 01:05:00 +0000