diff options
Diffstat (limited to 'flang/lib')
28 files changed, 958 insertions, 185 deletions
diff --git a/flang/lib/Evaluate/common.cpp b/flang/lib/Evaluate/common.cpp index 46c75a5..ed6a0ef 100644 --- a/flang/lib/Evaluate/common.cpp +++ b/flang/lib/Evaluate/common.cpp @@ -13,24 +13,28 @@ using namespace Fortran::parser::literals; namespace Fortran::evaluate { -void RealFlagWarnings( - FoldingContext &context, const RealFlags &flags, const char *operation) { +void FoldingContext::RealFlagWarnings( + const RealFlags &flags, const char *operation) { static constexpr auto warning{common::UsageWarning::FoldingException}; if (flags.test(RealFlag::Overflow)) { - context.Warn(warning, "overflow on %s"_warn_en_US, operation); + Warn(warning, "overflow on %s%s"_warn_en_US, operation, + realFlagWarningContext_); } if (flags.test(RealFlag::DivideByZero)) { if (std::strcmp(operation, "division") == 0) { - context.Warn(warning, "division by zero"_warn_en_US); + Warn(warning, "division by zero%s"_warn_en_US, realFlagWarningContext_); } else { - context.Warn(warning, "division by zero on %s"_warn_en_US, operation); + Warn(warning, "division by zero on %s%s"_warn_en_US, operation, + realFlagWarningContext_); } } if (flags.test(RealFlag::InvalidArgument)) { - context.Warn(warning, "invalid argument on %s"_warn_en_US, operation); + Warn(warning, "invalid argument on %s%s"_warn_en_US, operation, + realFlagWarningContext_); } if (flags.test(RealFlag::Underflow)) { - context.Warn(warning, "underflow on %s"_warn_en_US, operation); + Warn(warning, "underflow on %s%s"_warn_en_US, operation, + realFlagWarningContext_); } } diff --git a/flang/lib/Evaluate/fold-implementation.h b/flang/lib/Evaluate/fold-implementation.h index 3fdf3a6..52ea627 100644 --- a/flang/lib/Evaluate/fold-implementation.h +++ b/flang/lib/Evaluate/fold-implementation.h @@ -1862,7 +1862,7 @@ Expr<TO> FoldOperation( std::snprintf(buffer, sizeof buffer, "INTEGER(%d) to REAL(%d) conversion", Operand::kind, TO::kind); - RealFlagWarnings(ctx, converted.flags, buffer); + ctx.RealFlagWarnings(converted.flags, buffer); } return ScalarConstantToExpr(std::move(converted.value)); } else if constexpr (FromCat == TypeCategory::Real) { @@ -1871,7 +1871,7 @@ Expr<TO> FoldOperation( if (!converted.flags.empty()) { std::snprintf(buffer, sizeof buffer, "REAL(%d) to REAL(%d) conversion", Operand::kind, TO::kind); - RealFlagWarnings(ctx, converted.flags, buffer); + ctx.RealFlagWarnings(converted.flags, buffer); } if (ctx.targetCharacteristics().areSubnormalsFlushedToZero()) { converted.value = converted.value.FlushSubnormalToZero(); @@ -2012,7 +2012,7 @@ Expr<T> FoldOperation(FoldingContext &context, Add<T> &&x) { } else { auto sum{folded->first.Add( folded->second, context.targetCharacteristics().roundingMode())}; - RealFlagWarnings(context, sum.flags, "addition"); + context.RealFlagWarnings(sum.flags, "addition"); if (context.targetCharacteristics().areSubnormalsFlushedToZero()) { sum.value = sum.value.FlushSubnormalToZero(); } @@ -2041,7 +2041,7 @@ Expr<T> FoldOperation(FoldingContext &context, Subtract<T> &&x) { } else { auto difference{folded->first.Subtract( folded->second, context.targetCharacteristics().roundingMode())}; - RealFlagWarnings(context, difference.flags, "subtraction"); + context.RealFlagWarnings(difference.flags, "subtraction"); if (context.targetCharacteristics().areSubnormalsFlushedToZero()) { difference.value = difference.value.FlushSubnormalToZero(); } @@ -2070,7 +2070,7 @@ Expr<T> FoldOperation(FoldingContext &context, Multiply<T> &&x) { } else { auto product{folded->first.Multiply( folded->second, context.targetCharacteristics().roundingMode())}; - RealFlagWarnings(context, product.flags, "multiplication"); + context.RealFlagWarnings(product.flags, "multiplication"); if (context.targetCharacteristics().areSubnormalsFlushedToZero()) { product.value = product.value.FlushSubnormalToZero(); } @@ -2141,7 +2141,7 @@ Expr<T> FoldOperation(FoldingContext &context, Divide<T> &&x) { } } if (!isCanonicalNaNOrInf) { - RealFlagWarnings(context, quotient.flags, "division"); + context.RealFlagWarnings(quotient.flags, "division"); } if (context.targetCharacteristics().areSubnormalsFlushedToZero()) { quotient.value = quotient.value.FlushSubnormalToZero(); @@ -2201,7 +2201,7 @@ Expr<T> FoldOperation(FoldingContext &context, RealToIntPower<T> &&x) { [&](auto &y) -> Expr<T> { if (auto folded{OperandsAreConstants(x.left(), y)}) { auto power{evaluate::IntPower(folded->first, folded->second)}; - RealFlagWarnings(context, power.flags, "power with INTEGER exponent"); + context.RealFlagWarnings(power.flags, "power with INTEGER exponent"); if (context.targetCharacteristics().areSubnormalsFlushedToZero()) { power.value = power.value.FlushSubnormalToZero(); } diff --git a/flang/lib/Evaluate/host.cpp b/flang/lib/Evaluate/host.cpp index 25409ac..bf02496 100644 --- a/flang/lib/Evaluate/host.cpp +++ b/flang/lib/Evaluate/host.cpp @@ -140,8 +140,8 @@ void HostFloatingPointEnvironment::CheckAndRestoreFloatingPointEnvironment( } if (!flags_.empty()) { - RealFlagWarnings( - context, flags_, "evaluation of intrinsic function or operation"); + context.RealFlagWarnings( + flags_, "evaluation of intrinsic function or operation"); } errno = 0; if (fesetenv(&originalFenv_) != 0) { diff --git a/flang/lib/Evaluate/intrinsics-library.cpp b/flang/lib/Evaluate/intrinsics-library.cpp index 9820aa3..d8af524 100644 --- a/flang/lib/Evaluate/intrinsics-library.cpp +++ b/flang/lib/Evaluate/intrinsics-library.cpp @@ -1043,7 +1043,7 @@ std::optional<HostRuntimeWrapper> GetHostRuntimeWrapper(const std::string &name, if (const auto *hostFunction{ SearchHostRuntime(name, biggerResultType, biggerArgTypes)}) { auto hostFolderWithChecks{AddArgumentVerifierIfAny(name, *hostFunction)}; - return [hostFunction, resultType, hostFolderWithChecks]( + return [hostFunction, resultType, hostFolderWithChecks, name]( FoldingContext &context, std::vector<Expr<SomeType>> &&args) { auto nArgs{args.size()}; for (size_t i{0}; i < nArgs; ++i) { @@ -1051,6 +1051,8 @@ std::optional<HostRuntimeWrapper> GetHostRuntimeWrapper(const std::string &name, ConvertToType(hostFunction->argumentTypes[i], std::move(args[i])) .value()); } + auto restorer{context.SetRealFlagWarningContext( + " after folding a call to '"s + name + "'"s)}; return Fold(context, ConvertToType( resultType, hostFolderWithChecks(context, std::move(args))) diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp index 6e72987..0f4b39a 100644 --- a/flang/lib/Lower/Bridge.cpp +++ b/flang/lib/Lower/Bridge.cpp @@ -4876,6 +4876,10 @@ private: mlir::Value shape = builder->genShape(loc, lbounds, extents); rhsBox = fir::ReboxOp::create(*builder, loc, lhsBoxType, rhsBox, shape, /*slice=*/mlir::Value{}); + } else if (fir::isClassStarType(lhsBoxType) && + !fir::ConvertOp::canBeConverted(rhsBoxType, lhsBoxType)) { + rhsBox = fir::ReboxOp::create(*builder, loc, lhsBoxType, rhsBox, + mlir::Value{}, mlir::Value{}); } return rhsBox; } diff --git a/flang/lib/Lower/OpenMP/Clauses.cpp b/flang/lib/Lower/OpenMP/Clauses.cpp index d39f9dd..0f60b47 100644 --- a/flang/lib/Lower/OpenMP/Clauses.cpp +++ b/flang/lib/Lower/OpenMP/Clauses.cpp @@ -1482,6 +1482,21 @@ ThreadLimit make(const parser::OmpClause::ThreadLimit &inp, return ThreadLimit{/*Threadlim=*/makeExpr(inp.v, semaCtx)}; } +Threadset make(const parser::OmpClause::Threadset &inp, + semantics::SemanticsContext &semaCtx) { + // inp.v -> parser::OmpThreadsetClause + using wrapped = parser::OmpThreadsetClause; + + CLAUSET_ENUM_CONVERT( // + convert, wrapped::ThreadsetPolicy, Threadset::ThreadsetPolicy, + // clang-format off + MS(Omp_Pool, Omp_Pool) + MS(Omp_Team, Omp_Team) + // clang-format on + ); + return Threadset{/*ThreadsetPolicy=*/convert(inp.v.v)}; +} + // Threadprivate: empty // Threads: empty diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp index 39bac81..15ea845 100644 --- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp +++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp @@ -50,6 +50,7 @@ #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/LLVMIR/LLVMTypes.h" #include "mlir/Dialect/Math/IR/Math.h" +#include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" @@ -358,6 +359,14 @@ static constexpr IntrinsicHandler handlers[]{ &I::genBarrierInit, {{{"barrier", asAddr}, {"count", asValue}}}, /*isElemental=*/false}, + {"barrier_try_wait", + &I::genBarrierTryWait, + {{{"barrier", asAddr}, {"token", asValue}}}, + /*isElemental=*/false}, + {"barrier_try_wait_sleep", + &I::genBarrierTryWaitSleep, + {{{"barrier", asAddr}, {"token", asValue}, {"ns", asValue}}}, + /*isElemental=*/false}, {"bessel_jn", &I::genBesselJn, {{{"n1", asValue}, {"n2", asValue}, {"x", asValue}}}, @@ -1036,10 +1045,87 @@ static constexpr IntrinsicHandler handlers[]{ {"dst", asAddr}, {"nbytes", asValue}}}, /*isElemental=*/false}, + {"tma_bulk_ldc4", + &I::genTMABulkLoadC4, + {{{"barrier", asAddr}, + {"src", asAddr}, + {"dst", asAddr}, + {"nelems", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_ldc8", + &I::genTMABulkLoadC8, + {{{"barrier", asAddr}, + {"src", asAddr}, + {"dst", asAddr}, + {"nelems", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_ldi4", + &I::genTMABulkLoadI4, + {{{"barrier", asAddr}, + {"src", asAddr}, + {"dst", asAddr}, + {"nelems", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_ldi8", + &I::genTMABulkLoadI8, + {{{"barrier", asAddr}, + {"src", asAddr}, + {"dst", asAddr}, + {"nelems", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_ldr2", + &I::genTMABulkLoadR2, + {{{"barrier", asAddr}, + {"src", asAddr}, + {"dst", asAddr}, + {"nelems", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_ldr4", + &I::genTMABulkLoadR4, + {{{"barrier", asAddr}, + {"src", asAddr}, + {"dst", asAddr}, + {"nelems", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_ldr8", + &I::genTMABulkLoadR8, + {{{"barrier", asAddr}, + {"src", asAddr}, + {"dst", asAddr}, + {"nelems", asValue}}}, + /*isElemental=*/false}, {"tma_bulk_s2g", &I::genTMABulkS2G, {{{"src", asAddr}, {"dst", asAddr}, {"nbytes", asValue}}}, /*isElemental=*/false}, + {"tma_bulk_store_c4", + &I::genTMABulkStoreC4, + {{{"src", asAddr}, {"dst", asAddr}, {"count", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_store_c8", + &I::genTMABulkStoreC8, + {{{"src", asAddr}, {"dst", asAddr}, {"count", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_store_i4", + &I::genTMABulkStoreI4, + {{{"src", asAddr}, {"dst", asAddr}, {"count", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_store_i8", + &I::genTMABulkStoreI8, + {{{"src", asAddr}, {"dst", asAddr}, {"count", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_store_r2", + &I::genTMABulkStoreR2, + {{{"src", asAddr}, {"dst", asAddr}, {"count", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_store_r4", + &I::genTMABulkStoreR4, + {{{"src", asAddr}, {"dst", asAddr}, {"count", asValue}}}, + /*isElemental=*/false}, + {"tma_bulk_store_r8", + &I::genTMABulkStoreR8, + {{{"src", asAddr}, {"dst", asAddr}, {"count", asValue}}}, + /*isElemental=*/false}, {"tma_bulk_wait_group", &I::genTMABulkWaitGroup, {{}}, @@ -3273,8 +3359,8 @@ void IntrinsicLibrary::genBarrierInit(llvm::ArrayRef<fir::ExtendedValue> args) { assert(args.size() == 2); mlir::Value barrier = convertPtrToNVVMSpace( builder, loc, fir::getBase(args[0]), mlir::NVVM::NVVMMemorySpace::Shared); - mlir::NVVM::MBarrierInitSharedOp::create(builder, loc, barrier, - fir::getBase(args[1]), {}); + mlir::NVVM::MBarrierInitOp::create(builder, loc, barrier, + fir::getBase(args[1]), {}); auto kind = mlir::NVVM::ProxyKindAttr::get( builder.getContext(), mlir::NVVM::ProxyKind::async_shared); auto space = mlir::NVVM::SharedSpaceAttr::get( @@ -3282,6 +3368,57 @@ void IntrinsicLibrary::genBarrierInit(llvm::ArrayRef<fir::ExtendedValue> args) { mlir::NVVM::FenceProxyOp::create(builder, loc, kind, space); } +// BARRIER_TRY_WAIT (CUDA) +mlir::Value +IntrinsicLibrary::genBarrierTryWait(mlir::Type resultType, + llvm::ArrayRef<mlir::Value> args) { + assert(args.size() == 2); + mlir::Value res = fir::AllocaOp::create(builder, loc, resultType); + mlir::Value zero = builder.createIntegerConstant(loc, resultType, 0); + fir::StoreOp::create(builder, loc, zero, res); + mlir::Value ns = + builder.createIntegerConstant(loc, builder.getI32Type(), 1000000); + mlir::Value load = fir::LoadOp::create(builder, loc, res); + auto whileOp = mlir::scf::WhileOp::create( + builder, loc, mlir::TypeRange{resultType}, mlir::ValueRange{load}); + mlir::Block *beforeBlock = builder.createBlock(&whileOp.getBefore()); + mlir::Value beforeArg = beforeBlock->addArgument(resultType, loc); + builder.setInsertionPointToStart(beforeBlock); + mlir::Value condition = mlir::arith::CmpIOp::create( + builder, loc, mlir::arith::CmpIPredicate::ne, beforeArg, zero); + mlir::scf::ConditionOp::create(builder, loc, condition, beforeArg); + mlir::Block *afterBlock = builder.createBlock(&whileOp.getAfter()); + afterBlock->addArgument(resultType, loc); + builder.setInsertionPointToStart(afterBlock); + auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(builder.getContext()); + auto barrier = builder.createConvert(loc, llvmPtrTy, args[0]); + mlir::Value ret = + mlir::NVVM::InlinePtxOp::create( + builder, loc, {resultType}, {barrier, args[1], ns}, {}, + ".reg .pred p; mbarrier.try_wait.shared.b64 p, [%1], %2, %3; " + "selp.b32 %0, 1, 0, p;", + {}) + .getResult(0); + mlir::scf::YieldOp::create(builder, loc, ret); + builder.setInsertionPointAfter(whileOp); + return whileOp.getResult(0); +} + +// BARRIER_TRY_WAIT_SLEEP (CUDA) +mlir::Value +IntrinsicLibrary::genBarrierTryWaitSleep(mlir::Type resultType, + llvm::ArrayRef<mlir::Value> args) { + assert(args.size() == 3); + auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(builder.getContext()); + auto barrier = builder.createConvert(loc, llvmPtrTy, args[0]); + return mlir::NVVM::InlinePtxOp::create( + builder, loc, {resultType}, {barrier, args[1], args[2]}, {}, + ".reg .pred p; mbarrier.try_wait.shared.b64 p, [%1], %2, %3; " + "selp.b32 %0, 1, 0, p;", + {}) + .getResult(0); +} + // BESSEL_JN fir::ExtendedValue IntrinsicLibrary::genBesselJn(mlir::Type resultType, @@ -9218,6 +9355,95 @@ void IntrinsicLibrary::genTMABulkG2S(llvm::ArrayRef<fir::ExtendedValue> args) { builder, loc, dst, src, barrier, fir::getBase(args[3]), {}, {}); } +static void genTMABulkLoad(fir::FirOpBuilder &builder, mlir::Location loc, + mlir::Value barrier, mlir::Value src, + mlir::Value dst, mlir::Value nelem, + mlir::Value eleSize) { + 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); + dst = builder.createConvert(loc, llvmPtrTy, dst); + src = builder.createConvert(loc, llvmPtrTy, src); + mlir::NVVM::InlinePtxOp::create( + builder, loc, mlir::TypeRange{}, {dst, src, size, barrier}, {}, + "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], " + "[%1], %2, [%3];", + {}); + mlir::NVVM::InlinePtxOp::create( + builder, loc, mlir::TypeRange{}, {barrier, size}, {}, + "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;", {}); +} + +// TMA_BULK_LOADC4 +void IntrinsicLibrary::genTMABulkLoadC4( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 4); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 8); + genTMABulkLoad(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), fir::getBase(args[3]), eleSize); +} + +// TMA_BULK_LOADC8 +void IntrinsicLibrary::genTMABulkLoadC8( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 4); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 16); + genTMABulkLoad(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), fir::getBase(args[3]), eleSize); +} + +// TMA_BULK_LOADI4 +void IntrinsicLibrary::genTMABulkLoadI4( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 4); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 4); + genTMABulkLoad(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), fir::getBase(args[3]), eleSize); +} + +// TMA_BULK_LOADI8 +void IntrinsicLibrary::genTMABulkLoadI8( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 4); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 8); + genTMABulkLoad(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), fir::getBase(args[3]), eleSize); +} + +// TMA_BULK_LOADR2 +void IntrinsicLibrary::genTMABulkLoadR2( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 4); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 2); + genTMABulkLoad(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), fir::getBase(args[3]), eleSize); +} + +// TMA_BULK_LOADR4 +void IntrinsicLibrary::genTMABulkLoadR4( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 4); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 4); + genTMABulkLoad(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), fir::getBase(args[3]), eleSize); +} + +// TMA_BULK_LOADR8 +void IntrinsicLibrary::genTMABulkLoadR8( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 4); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 8); + genTMABulkLoad(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), fir::getBase(args[3]), eleSize); +} + // TMA_BULK_S2G (CUDA) void IntrinsicLibrary::genTMABulkS2G(llvm::ArrayRef<fir::ExtendedValue> args) { assert(args.size() == 3); @@ -9227,6 +9453,97 @@ void IntrinsicLibrary::genTMABulkS2G(llvm::ArrayRef<fir::ExtendedValue> args) { mlir::NVVM::NVVMMemorySpace::Global); mlir::NVVM::CpAsyncBulkSharedCTAToGlobalOp::create( builder, loc, dst, src, fir::getBase(args[2]), {}, {}); + + mlir::NVVM::InlinePtxOp::create(builder, loc, mlir::TypeRange{}, {}, {}, + "cp.async.bulk.commit_group", {}); + mlir::NVVM::CpAsyncBulkWaitGroupOp::create(builder, loc, + builder.getI32IntegerAttr(0), {}); +} + +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); + src = convertPtrToNVVMSpace(builder, loc, src, + mlir::NVVM::NVVMMemorySpace::Shared); + dst = convertPtrToNVVMSpace(builder, loc, dst, + mlir::NVVM::NVVMMemorySpace::Global); + mlir::NVVM::CpAsyncBulkSharedCTAToGlobalOp::create(builder, loc, dst, src, + size, {}, {}); + mlir::NVVM::InlinePtxOp::create(builder, loc, mlir::TypeRange{}, {}, {}, + "cp.async.bulk.commit_group", {}); + mlir::NVVM::CpAsyncBulkWaitGroupOp::create(builder, loc, + builder.getI32IntegerAttr(0), {}); +} + +// TMA_BULK_STORE_C4 (CUDA) +void IntrinsicLibrary::genTMABulkStoreC4( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 3); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 8); + genTMABulkStore(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), eleSize); +} + +// TMA_BULK_STORE_C8 (CUDA) +void IntrinsicLibrary::genTMABulkStoreC8( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 3); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 16); + genTMABulkStore(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), eleSize); +} + +// TMA_BULK_STORE_I4 (CUDA) +void IntrinsicLibrary::genTMABulkStoreI4( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 3); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 4); + genTMABulkStore(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), eleSize); +} + +// TMA_BULK_STORE_I8 (CUDA) +void IntrinsicLibrary::genTMABulkStoreI8( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 3); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 8); + genTMABulkStore(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), eleSize); +} + +// TMA_BULK_STORE_R2 (CUDA) +void IntrinsicLibrary::genTMABulkStoreR2( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 3); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 2); + genTMABulkStore(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), eleSize); +} + +// TMA_BULK_STORE_R4 (CUDA) +void IntrinsicLibrary::genTMABulkStoreR4( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 3); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 4); + genTMABulkStore(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), eleSize); +} + +// TMA_BULK_STORE_R8 (CUDA) +void IntrinsicLibrary::genTMABulkStoreR8( + llvm::ArrayRef<fir::ExtendedValue> args) { + assert(args.size() == 3); + mlir::Value eleSize = + builder.createIntegerConstant(loc, builder.getI32Type(), 8); + genTMABulkStore(builder, loc, fir::getBase(args[0]), fir::getBase(args[1]), + fir::getBase(args[2]), eleSize); } // TMA_BULK_WAIT_GROUP (CUDA) diff --git a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp index 0776346..8ca2869 100644 --- a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp +++ b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp @@ -143,7 +143,8 @@ public: llvm::SmallVector<mlir::Type> operandsTypes; for (auto arg : gpuLaunchFunc.getKernelOperands()) operandsTypes.push_back(arg.getType()); - auto fctTy = mlir::FunctionType::get(&context, operandsTypes, {}); + auto fctTy = mlir::FunctionType::get(&context, operandsTypes, + gpuLaunchFunc.getResultTypes()); if (!hasPortableSignature(fctTy, op)) convertCallOp(gpuLaunchFunc, fctTy); } else if (auto addr = mlir::dyn_cast<fir::AddrOfOp>(op)) { @@ -520,10 +521,14 @@ public: llvm::SmallVector<mlir::Value, 1> newCallResults; // TODO propagate/update call argument and result attributes. if constexpr (std::is_same_v<std::decay_t<A>, mlir::gpu::LaunchFuncOp>) { + mlir::Value asyncToken = callOp.getAsyncToken(); auto newCall = A::create(*rewriter, loc, callOp.getKernel(), callOp.getGridSizeOperandValues(), callOp.getBlockSizeOperandValues(), - callOp.getDynamicSharedMemorySize(), newOpers); + callOp.getDynamicSharedMemorySize(), newOpers, + asyncToken ? asyncToken.getType() : nullptr, + callOp.getAsyncDependencies(), + /*clusterSize=*/std::nullopt); if (callOp.getClusterSizeX()) newCall.getClusterSizeXMutable().assign(callOp.getClusterSizeX()); if (callOp.getClusterSizeY()) diff --git a/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt b/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt index ef67ab1..898fb00 100644 --- a/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt +++ b/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt @@ -2,6 +2,7 @@ get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) add_flang_library(FIROpenACCSupport FIROpenACCAttributes.cpp + FIROpenACCOpsInterfaces.cpp FIROpenACCTypeInterfaces.cpp RegisterOpenACCExtensions.cpp diff --git a/flang/lib/Optimizer/OpenACC/Support/FIROpenACCOpsInterfaces.cpp b/flang/lib/Optimizer/OpenACC/Support/FIROpenACCOpsInterfaces.cpp new file mode 100644 index 0000000..c1734be --- /dev/null +++ b/flang/lib/Optimizer/OpenACC/Support/FIROpenACCOpsInterfaces.cpp @@ -0,0 +1,62 @@ +//===-- FIROpenACCOpsInterfaces.cpp ---------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Implementation of external operation interfaces for FIR. +// +//===----------------------------------------------------------------------===// + +#include "flang/Optimizer/OpenACC/Support/FIROpenACCOpsInterfaces.h" + +#include "flang/Optimizer/Dialect/FIROps.h" +#include "flang/Optimizer/HLFIR/HLFIROps.h" + +namespace fir::acc { + +template <> +mlir::Value PartialEntityAccessModel<fir::ArrayCoorOp>::getBaseEntity( + mlir::Operation *op) const { + return mlir::cast<fir::ArrayCoorOp>(op).getMemref(); +} + +template <> +mlir::Value PartialEntityAccessModel<fir::CoordinateOp>::getBaseEntity( + mlir::Operation *op) const { + return mlir::cast<fir::CoordinateOp>(op).getRef(); +} + +template <> +mlir::Value PartialEntityAccessModel<hlfir::DesignateOp>::getBaseEntity( + mlir::Operation *op) const { + return mlir::cast<hlfir::DesignateOp>(op).getMemref(); +} + +mlir::Value PartialEntityAccessModel<fir::DeclareOp>::getBaseEntity( + mlir::Operation *op) const { + return mlir::cast<fir::DeclareOp>(op).getStorage(); +} + +bool PartialEntityAccessModel<fir::DeclareOp>::isCompleteView( + mlir::Operation *op) const { + // Return false (partial view) only if storage is present + // Return true (complete view) if storage is absent + return !getBaseEntity(op); +} + +mlir::Value PartialEntityAccessModel<hlfir::DeclareOp>::getBaseEntity( + mlir::Operation *op) const { + return mlir::cast<hlfir::DeclareOp>(op).getStorage(); +} + +bool PartialEntityAccessModel<hlfir::DeclareOp>::isCompleteView( + mlir::Operation *op) const { + // Return false (partial view) only if storage is present + // Return true (complete view) if storage is absent + return !getBaseEntity(op); +} + +} // namespace fir::acc diff --git a/flang/lib/Optimizer/OpenACC/Support/FIROpenACCTypeInterfaces.cpp b/flang/lib/Optimizer/OpenACC/Support/FIROpenACCTypeInterfaces.cpp index ed9e41c..ae0f5fb8 100644 --- a/flang/lib/Optimizer/OpenACC/Support/FIROpenACCTypeInterfaces.cpp +++ b/flang/lib/Optimizer/OpenACC/Support/FIROpenACCTypeInterfaces.cpp @@ -193,6 +193,28 @@ OpenACCMappableModel<fir::PointerType>::getOffsetInBytes( mlir::Type type, mlir::Value var, mlir::ValueRange accBounds, const mlir::DataLayout &dataLayout) const; +template <typename Ty> +bool OpenACCMappableModel<Ty>::hasUnknownDimensions(mlir::Type type) const { + assert(fir::isa_ref_type(type) && "expected FIR reference type"); + return fir::hasDynamicSize(fir::unwrapRefType(type)); +} + +template bool OpenACCMappableModel<fir::ReferenceType>::hasUnknownDimensions( + mlir::Type type) const; + +template bool OpenACCMappableModel<fir::HeapType>::hasUnknownDimensions( + mlir::Type type) const; + +template bool OpenACCMappableModel<fir::PointerType>::hasUnknownDimensions( + mlir::Type type) const; + +template <> +bool OpenACCMappableModel<fir::BaseBoxType>::hasUnknownDimensions( + mlir::Type type) const { + // Descriptor-based entities have dimensions encoded. + return false; +} + static llvm::SmallVector<mlir::Value> generateSeqTyAccBounds(fir::SequenceType seqType, mlir::Value var, mlir::OpBuilder &builder) { diff --git a/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp b/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp index 717bf34..d71c40d 100644 --- a/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp +++ b/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp @@ -11,8 +11,13 @@ //===----------------------------------------------------------------------===// #include "flang/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.h" + #include "flang/Optimizer/Dialect/FIRDialect.h" +#include "flang/Optimizer/Dialect/FIROps.h" #include "flang/Optimizer/Dialect/FIRType.h" +#include "flang/Optimizer/HLFIR/HLFIRDialect.h" +#include "flang/Optimizer/HLFIR/HLFIROps.h" +#include "flang/Optimizer/OpenACC/Support/FIROpenACCOpsInterfaces.h" #include "flang/Optimizer/OpenACC/Support/FIROpenACCTypeInterfaces.h" namespace fir::acc { @@ -37,7 +42,24 @@ void registerOpenACCExtensions(mlir::DialectRegistry ®istry) { fir::LLVMPointerType::attachInterface< OpenACCPointerLikeModel<fir::LLVMPointerType>>(*ctx); + + fir::ArrayCoorOp::attachInterface< + PartialEntityAccessModel<fir::ArrayCoorOp>>(*ctx); + fir::CoordinateOp::attachInterface< + PartialEntityAccessModel<fir::CoordinateOp>>(*ctx); + fir::DeclareOp::attachInterface<PartialEntityAccessModel<fir::DeclareOp>>( + *ctx); }); + + // Register HLFIR operation interfaces + registry.addExtension( + +[](mlir::MLIRContext *ctx, hlfir::hlfirDialect *dialect) { + hlfir::DesignateOp::attachInterface< + PartialEntityAccessModel<hlfir::DesignateOp>>(*ctx); + hlfir::DeclareOp::attachInterface< + PartialEntityAccessModel<hlfir::DeclareOp>>(*ctx); + }); + registerAttrsExtensions(registry); } diff --git a/flang/lib/Parser/openmp-parsers.cpp b/flang/lib/Parser/openmp-parsers.cpp index d1e081c..4159d2e 100644 --- a/flang/lib/Parser/openmp-parsers.cpp +++ b/flang/lib/Parser/openmp-parsers.cpp @@ -275,6 +275,13 @@ struct SpecificModifierParser { // --- Iterator helpers ----------------------------------------------- +static EntityDecl MakeEntityDecl(ObjectName &&name) { + return EntityDecl( + /*ObjectName=*/std::move(name), std::optional<ArraySpec>{}, + std::optional<CoarraySpec>{}, std::optional<CharLength>{}, + std::optional<Initialization>{}); +} + // [5.0:47:17-18] In an iterator-specifier, if the iterator-type is not // specified then the type of that iterator is default integer. // [5.0:49:14] The iterator-type must be an integer type. @@ -282,11 +289,7 @@ static std::list<EntityDecl> makeEntityList(std::list<ObjectName> &&names) { std::list<EntityDecl> entities; for (auto iter = names.begin(), end = names.end(); iter != end; ++iter) { - EntityDecl entityDecl( - /*ObjectName=*/std::move(*iter), std::optional<ArraySpec>{}, - std::optional<CoarraySpec>{}, std::optional<CharLength>{}, - std::optional<Initialization>{}); - entities.push_back(std::move(entityDecl)); + entities.push_back(MakeEntityDecl(std::move(*iter))); } return entities; } @@ -306,6 +309,217 @@ static TypeDeclarationStmt makeIterSpecDecl(std::list<ObjectName> &&names) { makeEntityList(std::move(names))); } +// --- Stylized expression handling ----------------------------------- + +// OpenMP has a concept of am "OpenMP stylized expression". Syntactially +// it looks like a typical Fortran expression (or statement), except: +// - the only variables allowed in it are OpenMP special variables, the +// exact set of these variables depends on the specific case of the +// stylized expression +// - the special OpenMP variables present may assume one or more types, +// and the expression should be semantically valid for each type. +// +// The stylized expression can be thought of as a template, which will be +// instantiated for each type provided somewhere in the context in which +// the stylized expression appears. +// +// AST nodes: +// - OmpStylizedExpression: contains the source string for the expression, +// plus the list of instances (OmpStylizedInstance). +// - OmpStylizedInstance: corresponds to the instantiation of the stylized +// expression for a specific type. The way that the type is specified is +// by creating declarations (OmpStylizedDeclaration) for the special +// variables. Together with the AST tree corresponding to the stylized +// expression the instantiation has enough information for semantic +// analysis. Each instance has its own scope, and the special variables +// have their own Symbol's (local to the scope). +// - OmpStylizedDeclaration: encapsulates the information that the visitors +// in resolve-names can use to "emulate" a declaration for a special +// variable and allow name resolution in the instantiation AST to work. +// +// Implementation specifics: +// The semantic analysis stores "evaluate::Expr" in each AST node rooted +// in parser::Expr (in the typedExpr member). The evaluate::Expr is specific +// to a given type, and so to allow different types for a given expression, +// for each type a separate copy of the parser::Expr subtree is created. +// Normally, AST nodes are non-copyable (copy-ctor is deleted), so to create +// several copies of a subtree, the same source string is parsed several +// times. The ParseState member in OmpStylizedExpression is the parser state +// immediately before the stylized expression. +// +// Initially, when OmpStylizedExpression is first created, the expression is +// parsed as if it was an actual code, but this parsing is only done to +// establish where the stylized expression ends (in the source). The source +// and the initial parser state are stored in the object, and the instance +// list is empty. +// Once the parsing of the containing OmpDirectiveSpecification completes, +// a post-processing "parser" (OmpStylizedInstanceCreator) executes. This +// post-processor examines the directive specification to see if it expects +// any stylized expressions to be contained in it, and then instantiates +// them for each such directive. + +template <typename A> struct NeverParser { + using resultType = A; + std::optional<resultType> Parse(ParseState &state) const { + // Always fail, but without any messages. + return std::nullopt; + } +}; + +template <typename A> constexpr auto never() { return NeverParser<A>{}; } + +// Parser for optional<T> which always succeeds and returns std::nullptr. +// It's only needed to produce "std::optional<CallStmt::Chevrons>" in +// CallStmt. +template <typename A, typename B = void> struct NullParser; +template <typename B> struct NullParser<std::optional<B>> { + using resultType = std::optional<B>; + std::optional<resultType> Parse(ParseState &) const { + return resultType{std::nullopt}; + } +}; + +template <typename A> constexpr auto null() { return NullParser<A>{}; } + +// OmpStylizedDeclaration and OmpStylizedInstance are helper classes, and +// don't correspond to anything in the source. Their parsers should still +// exist, but they should never be executed. +TYPE_PARSER(construct<OmpStylizedDeclaration>(never<OmpStylizedDeclaration>())) +TYPE_PARSER(construct<OmpStylizedInstance>(never<OmpStylizedInstance>())) + +TYPE_PARSER( // + construct<OmpStylizedInstance::Instance>(Parser<AssignmentStmt>{}) || + construct<OmpStylizedInstance::Instance>( + sourced(construct<CallStmt>(Parser<ProcedureDesignator>{}, + null<std::optional<CallStmt::Chevrons>>(), + parenthesized(optionalList(actualArgSpec))))) || + construct<OmpStylizedInstance::Instance>(indirect(expr))) + +struct OmpStylizedExpressionParser { + using resultType = OmpStylizedExpression; + + std::optional<resultType> Parse(ParseState &state) const { + auto *saved{new ParseState(state)}; + auto getSource{verbatim(Parser<OmpStylizedInstance::Instance>{} >> ok)}; + if (auto &&ok{getSource.Parse(state)}) { + OmpStylizedExpression result{std::list<OmpStylizedInstance>{}}; + result.source = ok->source; + result.state = saved; + // result.v remains empty + return std::move(result); + } + delete saved; + return std::nullopt; + } +}; + +static void Instantiate(OmpStylizedExpression &ose, + llvm::ArrayRef<const OmpTypeName *> types, llvm::ArrayRef<CharBlock> vars) { + // 1. For each var in the vars list, declare it with the corresponding + // type from types. + // 2. Run the parser to get the AST for the stylized expression. + // 3. Create OmpStylizedInstance and append it to the list in ose. + assert(types.size() == vars.size() && "List size mismatch"); + // A ParseState object is irreversibly modified during parsing (in + // particular, it cannot be rewound to an earlier position in the source). + // Because of that we need to create a local copy for each instantiation. + // If rewinding was possible, we could just use the current one, and we + // wouldn't need to save it in the AST node. + ParseState state{DEREF(ose.state)}; + + std::list<OmpStylizedDeclaration> decls; + for (auto [type, var] : llvm::zip_equal(types, vars)) { + decls.emplace_back(OmpStylizedDeclaration{ + common::Reference(*type), MakeEntityDecl(Name{var})}); + } + + if (auto &&instance{Parser<OmpStylizedInstance::Instance>{}.Parse(state)}) { + ose.v.emplace_back( + OmpStylizedInstance{std::move(decls), std::move(*instance)}); + } +} + +static void InstantiateForTypes(OmpStylizedExpression &ose, + const OmpTypeNameList &typeNames, llvm::ArrayRef<CharBlock> vars) { + // For each type in the type list, declare all variables in vars with + // that type, and complete the instantiation. + for (const OmpTypeName &t : typeNames.v) { + std::vector<const OmpTypeName *> types(vars.size(), &t); + Instantiate(ose, types, vars); + } +} + +static void InstantiateDeclareReduction(OmpDirectiveSpecification &spec) { + // There can be arguments/clauses that don't make sense, that analysis + // is left until semantic checks. Tolerate any unexpected stuff. + auto *rspec{GetFirstArgument<OmpReductionSpecifier>(spec)}; + if (!rspec) { + return; + } + + const OmpTypeNameList *typeNames{nullptr}; + + if (auto *cexpr{ + const_cast<OmpCombinerExpression *>(GetCombinerExpr(*rspec))}) { + typeNames = &std::get<OmpTypeNameList>(rspec->t); + + InstantiateForTypes(*cexpr, *typeNames, OmpCombinerExpression::Variables()); + delete cexpr->state; + cexpr->state = nullptr; + } else { + // If there are no types, there is nothing else to do. + return; + } + + for (const OmpClause &clause : spec.Clauses().v) { + llvm::omp::Clause id{clause.Id()}; + if (id == llvm::omp::Clause::OMPC_initializer) { + if (auto *iexpr{const_cast<OmpInitializerExpression *>( + GetInitializerExpr(clause))}) { + InstantiateForTypes( + *iexpr, *typeNames, OmpInitializerExpression::Variables()); + delete iexpr->state; + iexpr->state = nullptr; + } + } + } +} + +static void InstantiateStylizedDirective(OmpDirectiveSpecification &spec) { + const OmpDirectiveName &dirName{spec.DirName()}; + if (dirName.v == llvm::omp::Directive::OMPD_declare_reduction) { + InstantiateDeclareReduction(spec); + } +} + +template <typename P, + typename = std::enable_if_t< + std::is_same_v<typename P::resultType, OmpDirectiveSpecification>>> +struct OmpStylizedInstanceCreator { + using resultType = OmpDirectiveSpecification; + constexpr OmpStylizedInstanceCreator(P p) : parser_(p) {} + + std::optional<resultType> Parse(ParseState &state) const { + if (auto &&spec{parser_.Parse(state)}) { + InstantiateStylizedDirective(*spec); + return std::move(spec); + } + return std::nullopt; + } + +private: + const P parser_; +}; + +template <typename P> +OmpStylizedInstanceCreator(P) -> OmpStylizedInstanceCreator<P>; + +// --- Parsers for types ---------------------------------------------- + +TYPE_PARSER( // + sourced(construct<OmpTypeName>(Parser<DeclarationTypeSpec>{})) || + sourced(construct<OmpTypeName>(Parser<TypeSpec>{}))) + // --- Parsers for arguments ------------------------------------------ // At the moment these are only directive arguments. This is needed for @@ -366,10 +580,6 @@ struct OmpArgumentListParser { } }; -TYPE_PARSER( // - construct<OmpTypeName>(Parser<DeclarationTypeSpec>{}) || - construct<OmpTypeName>(Parser<TypeSpec>{})) - // 2.15.3.6 REDUCTION (reduction-identifier: variable-name-list) TYPE_PARSER(construct<OmpReductionIdentifier>(Parser<DefinedOperator>{}) || construct<OmpReductionIdentifier>(Parser<ProcedureDesignator>{})) @@ -1065,7 +1275,8 @@ TYPE_PARSER(construct<OmpOtherwiseClause>( TYPE_PARSER(construct<OmpWhenClause>( maybe(nonemptyList(Parser<OmpWhenClause::Modifier>{}) / ":"), - maybe(indirect(Parser<OmpDirectiveSpecification>{})))) + maybe(indirect( + OmpStylizedInstanceCreator(Parser<OmpDirectiveSpecification>{}))))) // OMP 5.2 12.6.1 grainsize([ prescriptiveness :] scalar-integer-expression) TYPE_PARSER(construct<OmpGrainsizeClause>( @@ -1777,12 +1988,7 @@ TYPE_PARSER( Parser<OpenMPInteropConstruct>{})) / endOfLine) -TYPE_PARSER(construct<OmpInitializerProc>(Parser<ProcedureDesignator>{}, - parenthesized(many(maybe(","_tok) >> Parser<ActualArgSpec>{})))) - -TYPE_PARSER(construct<OmpInitializerClause>( - construct<OmpInitializerClause>(assignmentStmt) || - construct<OmpInitializerClause>(Parser<OmpInitializerProc>{}))) +TYPE_PARSER(construct<OmpInitializerClause>(Parser<OmpInitializerExpression>{})) // OpenMP 5.2: 7.5.4 Declare Variant directive TYPE_PARSER(sourced(construct<OmpDeclareVariantDirective>( @@ -1794,7 +2000,7 @@ TYPE_PARSER(sourced(construct<OmpDeclareVariantDirective>( TYPE_PARSER(sourced(construct<OpenMPDeclareReductionConstruct>( predicated(Parser<OmpDirectiveName>{}, IsDirective(llvm::omp::Directive::OMPD_declare_reduction)) >= - Parser<OmpDirectiveSpecification>{}))) + OmpStylizedInstanceCreator(Parser<OmpDirectiveSpecification>{})))) // 2.10.6 Declare Target Construct TYPE_PARSER(sourced(construct<OpenMPDeclareTargetConstruct>( @@ -1832,8 +2038,8 @@ TYPE_PARSER(sourced(construct<OpenMPDeclareMapperConstruct>( IsDirective(llvm::omp::Directive::OMPD_declare_mapper)) >= Parser<OmpDirectiveSpecification>{}))) -TYPE_PARSER(construct<OmpCombinerExpression>(Parser<AssignmentStmt>{}) || - construct<OmpCombinerExpression>(Parser<FunctionReference>{})) +TYPE_PARSER(construct<OmpCombinerExpression>(OmpStylizedExpressionParser{})) +TYPE_PARSER(construct<OmpInitializerExpression>(OmpStylizedExpressionParser{})) TYPE_PARSER(sourced(construct<OpenMPCriticalConstruct>( OmpBlockConstructParser{llvm::omp::Directive::OMPD_critical}))) diff --git a/flang/lib/Parser/openmp-utils.cpp b/flang/lib/Parser/openmp-utils.cpp index 937a17f..95ad3f6 100644 --- a/flang/lib/Parser/openmp-utils.cpp +++ b/flang/lib/Parser/openmp-utils.cpp @@ -74,4 +74,16 @@ const BlockConstruct *GetFortranBlockConstruct( return nullptr; } +const OmpCombinerExpression *GetCombinerExpr( + const OmpReductionSpecifier &rspec) { + return addr_if(std::get<std::optional<OmpCombinerExpression>>(rspec.t)); +} + +const OmpInitializerExpression *GetInitializerExpr(const OmpClause &init) { + if (auto *wrapped{std::get_if<OmpClause::Initializer>(&init.u)}) { + return &wrapped->v.v; + } + return nullptr; +} + } // namespace Fortran::parser::omp diff --git a/flang/lib/Parser/parse-tree.cpp b/flang/lib/Parser/parse-tree.cpp index 8cbaa39..ad0016e 100644 --- a/flang/lib/Parser/parse-tree.cpp +++ b/flang/lib/Parser/parse-tree.cpp @@ -11,6 +11,7 @@ #include "flang/Common/indirection.h" #include "flang/Parser/tools.h" #include "flang/Parser/user-state.h" +#include "llvm/ADT/ArrayRef.h" #include "llvm/Frontend/OpenMP/OMP.h" #include "llvm/Support/raw_ostream.h" #include <algorithm> @@ -430,4 +431,30 @@ const OmpClauseList &OmpDirectiveSpecification::Clauses() const { } return empty; } + +static bool InitCharBlocksFromStrings(llvm::MutableArrayRef<CharBlock> blocks, + llvm::ArrayRef<std::string> strings) { + for (auto [i, n] : llvm::enumerate(strings)) { + blocks[i] = CharBlock(n); + } + return true; +} + +// The names should have static storage duration. Keep these names +// in a sigle place. +llvm::ArrayRef<CharBlock> OmpCombinerExpression::Variables() { + static std::string names[]{"omp_in", "omp_out"}; + static CharBlock vars[std::size(names)]; + + [[maybe_unused]] static bool init = InitCharBlocksFromStrings(vars, names); + return vars; +} + +llvm::ArrayRef<CharBlock> OmpInitializerExpression::Variables() { + static std::string names[]{"omp_orig", "omp_priv"}; + static CharBlock vars[std::size(names)]; + + [[maybe_unused]] static bool init = InitCharBlocksFromStrings(vars, names); + return vars; +} } // namespace Fortran::parser diff --git a/flang/lib/Parser/prescan.cpp b/flang/lib/Parser/prescan.cpp index 4739da0..efce8fc 100644 --- a/flang/lib/Parser/prescan.cpp +++ b/flang/lib/Parser/prescan.cpp @@ -557,7 +557,7 @@ bool Prescanner::MustSkipToEndOfLine() const { return true; // skip over ignored columns in right margin (73:80) } else if (*at_ == '!' && !inCharLiteral_ && (!inFixedForm_ || tabInCurrentLine_ || column_ != 6)) { - return !IsCompilerDirectiveSentinel(at_); + return !IsCompilerDirectiveSentinel(at_ + 1); } else { return false; } @@ -1642,6 +1642,17 @@ Prescanner::IsFixedFormCompilerDirectiveLine(const char *start) const { // This is a Continuation line, not an initial directive line. return std::nullopt; } + ++column, ++p; + } + if (isOpenMPConditional) { + for (; column <= fixedFormColumnLimit_; ++column, ++p) { + if (IsSpaceOrTab(p)) { + } else if (*p == '!') { + return std::nullopt; // !$ ! is a comment, not a directive + } else { + break; + } + } } if (const char *ss{IsCompilerDirectiveSentinel( sentinel, static_cast<std::size_t>(sp - sentinel))}) { @@ -1657,8 +1668,17 @@ Prescanner::IsFreeFormCompilerDirectiveLine(const char *start) const { p && *p++ == '!') { if (auto maybePair{IsCompilerDirectiveSentinel(p)}) { auto offset{static_cast<std::size_t>(p - start - 1)}; - return {LineClassification{LineClassification::Kind::CompilerDirective, - offset, maybePair->first}}; + const char *sentinel{maybePair->first}; + if ((sentinel[0] == '$' && sentinel[1] == '\0') || sentinel[1] == '@') { + if (const char *comment{IsFreeFormComment(maybePair->second)}) { + if (*comment == '!') { + // Conditional line comment - treat as comment + return std::nullopt; + } + } + } + return {LineClassification{ + LineClassification::Kind::CompilerDirective, offset, sentinel}}; } } return std::nullopt; diff --git a/flang/lib/Parser/unparse.cpp b/flang/lib/Parser/unparse.cpp index 20a8d2a..9b38cfc 100644 --- a/flang/lib/Parser/unparse.cpp +++ b/flang/lib/Parser/unparse.cpp @@ -2095,15 +2095,13 @@ public: // OpenMP Clauses & Directives void Unparse(const OmpArgumentList &x) { Walk(x.v, ", "); } + void Unparse(const OmpTypeNameList &x) { Walk(x.v, ", "); } void Unparse(const OmpBaseVariantNames &x) { Walk(std::get<0>(x.t)); // OmpObject Put(":"); Walk(std::get<1>(x.t)); // OmpObject } - void Unparse(const OmpTypeNameList &x) { // - Walk(x.v, ","); - } void Unparse(const OmpMapperSpecifier &x) { const auto &mapperName{std::get<std::string>(x.t)}; if (mapperName.find(llvm::omp::OmpDefaultMapperName) == std::string::npos) { @@ -2202,6 +2200,15 @@ public: unsigned ompVersion{langOpts_.OpenMPVersion}; Word(llvm::omp::getOpenMPDirectiveName(x.v, ompVersion)); } + void Unparse(const OmpStylizedDeclaration &x) { + // empty + } + void Unparse(const OmpStylizedExpression &x) { // + Put(x.source.ToString()); + } + void Unparse(const OmpStylizedInstance &x) { + // empty + } void Unparse(const OmpIteratorSpecifier &x) { Walk(std::get<TypeDeclarationStmt>(x.t)); Put(" = "); @@ -2511,29 +2518,11 @@ public: void Unparse(const OpenMPCriticalConstruct &x) { Unparse(static_cast<const OmpBlockConstruct &>(x)); } - void Unparse(const OmpInitializerProc &x) { - Walk(std::get<ProcedureDesignator>(x.t)); - Put("("); - Walk(std::get<std::list<ActualArgSpec>>(x.t)); - Put(")"); - } - void Unparse(const OmpInitializerClause &x) { - // Don't let the visitor go to the normal AssignmentStmt Unparse function, - // it adds an extra newline that we don't want. - if (const auto *assignment{std::get_if<AssignmentStmt>(&x.u)}) { - Walk(assignment->t, " = "); - } else { - Walk(x.u); - } + void Unparse(const OmpInitializerExpression &x) { + Unparse(static_cast<const OmpStylizedExpression &>(x)); } void Unparse(const OmpCombinerExpression &x) { - // Don't let the visitor go to the normal AssignmentStmt Unparse function, - // it adds an extra newline that we don't want. - if (const auto *assignment{std::get_if<AssignmentStmt>(&x.u)}) { - Walk(assignment->t, " = "); - } else { - Walk(x.u); - } + Unparse(static_cast<const OmpStylizedExpression &>(x)); } void Unparse(const OpenMPDeclareReductionConstruct &x) { BeginOpenMP(); diff --git a/flang/lib/Semantics/check-allocate.cpp b/flang/lib/Semantics/check-allocate.cpp index e019bbd..a411e20 100644 --- a/flang/lib/Semantics/check-allocate.cpp +++ b/flang/lib/Semantics/check-allocate.cpp @@ -26,6 +26,10 @@ struct AllocateCheckerInfo { std::optional<evaluate::DynamicType> sourceExprType; std::optional<parser::CharBlock> sourceExprLoc; std::optional<parser::CharBlock> typeSpecLoc; + std::optional<parser::CharBlock> statSource; + std::optional<parser::CharBlock> msgSource; + const SomeExpr *statVar{nullptr}; + const SomeExpr *msgVar{nullptr}; int sourceExprRank{0}; // only valid if gotMold || gotSource bool gotStat{false}; bool gotMsg{false}; @@ -141,12 +145,15 @@ static std::optional<AllocateCheckerInfo> CheckAllocateOptions( [&](const parser::StatOrErrmsg &statOrErr) { common::visit( common::visitors{ - [&](const parser::StatVariable &) { + [&](const parser::StatVariable &var) { if (info.gotStat) { // C943 context.Say( "STAT may not be duplicated in a ALLOCATE statement"_err_en_US); } info.gotStat = true; + info.statVar = GetExpr(context, var); + info.statSource = + parser::Unwrap<parser::Variable>(var)->GetSource(); }, [&](const parser::MsgVariable &var) { WarnOnDeferredLengthCharacterScalar(context, @@ -159,6 +166,9 @@ static std::optional<AllocateCheckerInfo> CheckAllocateOptions( "ERRMSG may not be duplicated in a ALLOCATE statement"_err_en_US); } info.gotMsg = true; + info.msgVar = GetExpr(context, var); + info.msgSource = + parser::Unwrap<parser::Variable>(var)->GetSource(); }, }, statOrErr.u); @@ -460,6 +470,16 @@ static bool HaveCompatibleLengths( } } +bool AreSameAllocation(const SomeExpr *root, const SomeExpr *path) { + if (root && path) { + // For now we just use equality of expressions. If we implement a more + // sophisticated alias analysis we should use it here. + return *root == *path; + } else { + return false; + } +} + bool AllocationCheckerHelper::RunChecks(SemanticsContext &context) { if (!ultimate_) { CHECK(context.AnyFatalError()); @@ -690,6 +710,17 @@ bool AllocationCheckerHelper::RunChecks(SemanticsContext &context) { "Object in ALLOCATE must have DEVICE attribute when STREAM option is specified"_err_en_US); } } + + if (const SomeExpr *allocObj{GetExpr(context, allocateObject_)}) { + if (AreSameAllocation(allocObj, allocateInfo_.statVar)) { + context.Say(allocateInfo_.statSource.value_or(name_.source), + "STAT variable in ALLOCATE must not be the variable being allocated"_err_en_US); + } + if (AreSameAllocation(allocObj, allocateInfo_.msgVar)) { + context.Say(allocateInfo_.msgSource.value_or(name_.source), + "ERRMSG variable in ALLOCATE must not be the variable being allocated"_err_en_US); + } + } return RunCoarrayRelatedChecks(context); } diff --git a/flang/lib/Semantics/check-allocate.h b/flang/lib/Semantics/check-allocate.h index e3f7f07..54f7380 100644 --- a/flang/lib/Semantics/check-allocate.h +++ b/flang/lib/Semantics/check-allocate.h @@ -24,5 +24,6 @@ public: private: SemanticsContext &context_; }; +bool AreSameAllocation(const SomeExpr *root, const SomeExpr *path); } // namespace Fortran::semantics #endif // FORTRAN_SEMANTICS_CHECK_ALLOCATE_H_ diff --git a/flang/lib/Semantics/check-call.cpp b/flang/lib/Semantics/check-call.cpp index c51d40b..995deaa 100644 --- a/flang/lib/Semantics/check-call.cpp +++ b/flang/lib/Semantics/check-call.cpp @@ -914,7 +914,8 @@ static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy, dummyName); } // INTENT(OUT) and INTENT(IN OUT) cases are caught elsewhere - } else { + } else if (!actualIsAllocatable && + !dummy.ignoreTKR.test(common::IgnoreTKR::Pointer)) { messages.Say( "ALLOCATABLE %s must be associated with an ALLOCATABLE actual argument"_err_en_US, dummyName); @@ -929,7 +930,8 @@ static void CheckExplicitDataArg(const characteristics::DummyDataObject &dummy, dummy, actual, *scope, /*isAssumedRank=*/dummyIsAssumedRank, actualIsPointer); } - } else if (!actualIsPointer) { + } else if (!actualIsPointer && + !dummy.ignoreTKR.test(common::IgnoreTKR::Pointer)) { messages.Say( "Actual argument associated with POINTER %s must also be POINTER unless INTENT(IN)"_err_en_US, dummyName); diff --git a/flang/lib/Semantics/check-deallocate.cpp b/flang/lib/Semantics/check-deallocate.cpp index c1ebc5f..e6ce1b3 100644 --- a/flang/lib/Semantics/check-deallocate.cpp +++ b/flang/lib/Semantics/check-deallocate.cpp @@ -7,51 +7,87 @@ //===----------------------------------------------------------------------===// #include "check-deallocate.h" +#include "check-allocate.h" #include "definable.h" #include "flang/Evaluate/type.h" #include "flang/Parser/message.h" #include "flang/Parser/parse-tree.h" #include "flang/Semantics/expression.h" #include "flang/Semantics/tools.h" +#include <optional> namespace Fortran::semantics { void DeallocateChecker::Leave(const parser::DeallocateStmt &deallocateStmt) { + bool gotStat{false}, gotMsg{false}; + const SomeExpr *statVar{nullptr}, *msgVar{nullptr}; + std::optional<parser::CharBlock> statSource; + std::optional<parser::CharBlock> msgSource; + for (const parser::StatOrErrmsg &deallocOpt : + std::get<std::list<parser::StatOrErrmsg>>(deallocateStmt.t)) { + common::visit( + common::visitors{ + [&](const parser::StatVariable &var) { + if (gotStat) { + context_.Say( + "STAT may not be duplicated in a DEALLOCATE statement"_err_en_US); + } + gotStat = true; + statVar = GetExpr(context_, var); + statSource = parser::Unwrap<parser::Variable>(var)->GetSource(); + }, + [&](const parser::MsgVariable &var) { + WarnOnDeferredLengthCharacterScalar(context_, + GetExpr(context_, var), + parser::UnwrapRef<parser::Variable>(var).GetSource(), + "ERRMSG="); + if (gotMsg) { + context_.Say( + "ERRMSG may not be duplicated in a DEALLOCATE statement"_err_en_US); + } + gotMsg = true; + msgVar = GetExpr(context_, var); + msgSource = parser::Unwrap<parser::Variable>(var)->GetSource(); + }, + }, + deallocOpt.u); + } for (const parser::AllocateObject &allocateObject : std::get<std::list<parser::AllocateObject>>(deallocateStmt.t)) { + parser::CharBlock source; common::visit( common::visitors{ [&](const parser::Name &name) { const Symbol *symbol{ name.symbol ? &name.symbol->GetUltimate() : nullptr}; - ; + source = name.source; if (context_.HasError(symbol)) { // already reported an error } else if (!IsVariableName(*symbol)) { - context_.Say(name.source, + context_.Say(source, "Name in DEALLOCATE statement must be a variable name"_err_en_US); } else if (!IsAllocatableOrObjectPointer(symbol)) { // C936 - context_.Say(name.source, + context_.Say(source, "Name in DEALLOCATE statement must have the ALLOCATABLE or POINTER attribute"_err_en_US); - } else if (auto whyNot{WhyNotDefinable(name.source, - context_.FindScope(name.source), - {DefinabilityFlag::PointerDefinition, - DefinabilityFlag::AcceptAllocatable, - DefinabilityFlag::PotentialDeallocation}, - *symbol)}) { + } else if (auto whyNot{ + WhyNotDefinable(source, context_.FindScope(source), + {DefinabilityFlag::PointerDefinition, + DefinabilityFlag::AcceptAllocatable, + DefinabilityFlag::PotentialDeallocation}, + *symbol)}) { // Catch problems with non-definability of the // pointer/allocatable context_ - .Say(name.source, + .Say(source, "Name in DEALLOCATE statement is not definable"_err_en_US) .Attach(std::move( whyNot->set_severity(parser::Severity::Because))); - } else if (auto whyNot{WhyNotDefinable(name.source, - context_.FindScope(name.source), - DefinabilityFlags{}, *symbol)}) { + } else if (auto whyNot{ + WhyNotDefinable(source, context_.FindScope(source), + DefinabilityFlags{}, *symbol)}) { // Catch problems with non-definability of the dynamic object context_ - .Say(name.source, + .Say(source, "Object in DEALLOCATE statement is not deallocatable"_err_en_US) .Attach(std::move( whyNot->set_severity(parser::Severity::Because))); @@ -62,13 +98,12 @@ void DeallocateChecker::Leave(const parser::DeallocateStmt &deallocateStmt) { [&](const parser::StructureComponent &structureComponent) { // Only perform structureComponent checks if it was successfully // analyzed by expression analysis. - auto source{structureComponent.component.source}; + source = structureComponent.component.source; if (const auto *expr{GetExpr(context_, allocateObject)}) { - if (const Symbol * - symbol{structureComponent.component.symbol - ? &structureComponent.component.symbol - ->GetUltimate() - : nullptr}; + if (const Symbol *symbol{structureComponent.component.symbol + ? &structureComponent.component.symbol + ->GetUltimate() + : nullptr}; !IsAllocatableOrObjectPointer(symbol)) { // F'2023 C936 context_.Say(source, "Component in DEALLOCATE statement must have the ALLOCATABLE or POINTER attribute"_err_en_US); @@ -99,32 +134,16 @@ void DeallocateChecker::Leave(const parser::DeallocateStmt &deallocateStmt) { }, }, allocateObject.u); - } - bool gotStat{false}, gotMsg{false}; - for (const parser::StatOrErrmsg &deallocOpt : - std::get<std::list<parser::StatOrErrmsg>>(deallocateStmt.t)) { - common::visit( - common::visitors{ - [&](const parser::StatVariable &) { - if (gotStat) { - context_.Say( - "STAT may not be duplicated in a DEALLOCATE statement"_err_en_US); - } - gotStat = true; - }, - [&](const parser::MsgVariable &var) { - WarnOnDeferredLengthCharacterScalar(context_, - GetExpr(context_, var), - parser::UnwrapRef<parser::Variable>(var).GetSource(), - "ERRMSG="); - if (gotMsg) { - context_.Say( - "ERRMSG may not be duplicated in a DEALLOCATE statement"_err_en_US); - } - gotMsg = true; - }, - }, - deallocOpt.u); + if (const SomeExpr *allocObj{GetExpr(context_, allocateObject)}) { + if (AreSameAllocation(allocObj, statVar)) { + context_.Say(statSource.value_or(source), + "STAT variable in DEALLOCATE must not be the variable being deallocated"_err_en_US); + } + if (AreSameAllocation(allocObj, msgVar)) { + context_.Say(msgSource.value_or(source), + "ERRMSG variable in DEALLOCATE must not be the variable being deallocated"_err_en_US); + } + } } } diff --git a/flang/lib/Semantics/check-declarations.cpp b/flang/lib/Semantics/check-declarations.cpp index 549ee83..de407d3 100644 --- a/flang/lib/Semantics/check-declarations.cpp +++ b/flang/lib/Semantics/check-declarations.cpp @@ -949,7 +949,8 @@ void CheckHelper::CheckObjectEntity( "!DIR$ IGNORE_TKR(R) may not apply in an ELEMENTAL procedure"_err_en_US); } if (IsPassedViaDescriptor(symbol)) { - if (IsAllocatableOrObjectPointer(&symbol)) { + if (IsAllocatableOrObjectPointer(&symbol) && + !ignoreTKR.test(common::IgnoreTKR::Pointer)) { if (inExplicitExternalInterface) { Warn(common::UsageWarning::IgnoreTKRUsage, "!DIR$ IGNORE_TKR should not apply to an allocatable or pointer"_warn_en_US); diff --git a/flang/lib/Semantics/check-omp-structure.cpp b/flang/lib/Semantics/check-omp-structure.cpp index e094458f..aaaf1ec 100644 --- a/flang/lib/Semantics/check-omp-structure.cpp +++ b/flang/lib/Semantics/check-omp-structure.cpp @@ -3390,6 +3390,7 @@ CHECK_SIMPLE_CLAUSE(Read, OMPC_read) CHECK_SIMPLE_CLAUSE(Threadprivate, OMPC_threadprivate) CHECK_SIMPLE_CLAUSE(Groupprivate, OMPC_groupprivate) CHECK_SIMPLE_CLAUSE(Threads, OMPC_threads) +CHECK_SIMPLE_CLAUSE(Threadset, OMPC_threadset) CHECK_SIMPLE_CLAUSE(Inbranch, OMPC_inbranch) CHECK_SIMPLE_CLAUSE(Link, OMPC_link) CHECK_SIMPLE_CLAUSE(Indirect, OMPC_indirect) diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp index 32aa6b1..c8167fd 100644 --- a/flang/lib/Semantics/expression.cpp +++ b/flang/lib/Semantics/expression.cpp @@ -834,7 +834,7 @@ Constant<TYPE> ReadRealLiteral( auto valWithFlags{ Scalar<TYPE>::Read(p, context.targetCharacteristics().roundingMode())}; CHECK(p == source.end()); - RealFlagWarnings(context, valWithFlags.flags, "conversion of REAL literal"); + context.RealFlagWarnings(valWithFlags.flags, "conversion of REAL literal"); auto value{valWithFlags.value}; if (context.targetCharacteristics().areSubnormalsFlushedToZero()) { value = value.FlushSubnormalToZero(); diff --git a/flang/lib/Semantics/mod-file.cpp b/flang/lib/Semantics/mod-file.cpp index 556259d..b419864 100644 --- a/flang/lib/Semantics/mod-file.cpp +++ b/flang/lib/Semantics/mod-file.cpp @@ -1021,6 +1021,9 @@ void ModFileWriter::PutObjectEntity( case common::IgnoreTKR::Contiguous: os << 'c'; break; + case common::IgnoreTKR::Pointer: + os << 'p'; + break; } }); os << ") " << symbol.name() << '\n'; diff --git a/flang/lib/Semantics/resolve-directives.cpp b/flang/lib/Semantics/resolve-directives.cpp index 196755e..628068f 100644 --- a/flang/lib/Semantics/resolve-directives.cpp +++ b/flang/lib/Semantics/resolve-directives.cpp @@ -26,6 +26,8 @@ #include "flang/Semantics/symbol.h" #include "flang/Semantics/tools.h" #include "flang/Support/Flags.h" +#include "llvm/ADT/StringMap.h" +#include "llvm/ADT/StringRef.h" #include "llvm/Frontend/OpenMP/OMP.h.inc" #include "llvm/Support/Debug.h" #include <list> @@ -453,6 +455,21 @@ public: return true; } + bool Pre(const parser::OmpStylizedDeclaration &x) { + static llvm::StringMap<Symbol::Flag> map{ + {"omp_in", Symbol::Flag::OmpInVar}, + {"omp_orig", Symbol::Flag::OmpOrigVar}, + {"omp_out", Symbol::Flag::OmpOutVar}, + {"omp_priv", Symbol::Flag::OmpPrivVar}, + }; + if (auto &name{std::get<parser::ObjectName>(x.var.t)}; name.symbol) { + if (auto found{map.find(name.ToString())}; found != map.end()) { + ResolveOmp(name, found->second, + const_cast<Scope &>(DEREF(name.symbol).owner())); + } + } + return false; + } bool Pre(const parser::OmpMetadirectiveDirective &x) { PushContext(x.v.source, llvm::omp::Directive::OMPD_metadirective); return true; diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp index 93faba7..220f1c9 100644 --- a/flang/lib/Semantics/resolve-names.cpp +++ b/flang/lib/Semantics/resolve-names.cpp @@ -1605,6 +1605,12 @@ public: Post(static_cast<const parser::OmpDirectiveSpecification &>(x)); } + void Post(const parser::OmpTypeName &); + bool Pre(const parser::OmpStylizedDeclaration &); + void Post(const parser::OmpStylizedDeclaration &); + bool Pre(const parser::OmpStylizedInstance &); + void Post(const parser::OmpStylizedInstance &); + bool Pre(const parser::OpenMPDeclareMapperConstruct &x) { AddOmpSourceRange(x.source); return true; @@ -1615,18 +1621,6 @@ public: return true; } - bool Pre(const parser::OmpInitializerProc &x) { - auto &procDes = std::get<parser::ProcedureDesignator>(x.t); - auto &name = std::get<parser::Name>(procDes.u); - auto *symbol{FindSymbol(NonDerivedTypeScope(), name)}; - if (!symbol) { - context().Say(name.source, - "Implicit subroutine declaration '%s' in DECLARE REDUCTION"_err_en_US, - name.source); - } - return true; - } - bool Pre(const parser::OmpDeclareVariantDirective &x) { AddOmpSourceRange(x.source); return true; @@ -1772,14 +1766,6 @@ public: messageHandler().set_currStmtSource(std::nullopt); } - bool Pre(const parser::OmpTypeName &x) { - BeginDeclTypeSpec(); - return true; - } - void Post(const parser::OmpTypeName &x) { // - EndDeclTypeSpec(); - } - bool Pre(const parser::OpenMPConstruct &x) { // Indicate that the current directive is not a declarative one. declaratives_.push_back(nullptr); @@ -1835,6 +1821,30 @@ void OmpVisitor::Post(const parser::OmpBlockConstruct &x) { } } +void OmpVisitor::Post(const parser::OmpTypeName &x) { + x.declTypeSpec = GetDeclTypeSpec(); +} + +bool OmpVisitor::Pre(const parser::OmpStylizedDeclaration &x) { + BeginDecl(); + Walk(x.type.get()); + Walk(x.var); + return true; +} + +void OmpVisitor::Post(const parser::OmpStylizedDeclaration &x) { // + EndDecl(); +} + +bool OmpVisitor::Pre(const parser::OmpStylizedInstance &x) { + PushScope(Scope::Kind::OtherConstruct, nullptr); + return true; +} + +void OmpVisitor::Post(const parser::OmpStylizedInstance &x) { // + PopScope(); +} + bool OmpVisitor::Pre(const parser::OmpMapClause &x) { auto &mods{OmpGetModifiers(x)}; if (auto *mapper{OmpGetUniqueModifier<parser::OmpMapper>(mods)}) { @@ -1969,51 +1979,20 @@ void OmpVisitor::ProcessReductionSpecifier( } } - auto &typeList{std::get<parser::OmpTypeNameList>(spec.t)}; - - // Create a temporary variable declaration for the four variables - // used in the reduction specifier and initializer (omp_out, omp_in, - // omp_priv and omp_orig), with the type in the typeList. - // - // In theory it would be possible to create only variables that are - // actually used, but that requires walking the entire parse-tree of the - // expressions, and finding the relevant variables [there may well be other - // variables involved too]. - // - // This allows doing semantic analysis where the type is a derived type - // e.g omp_out%x = omp_out%x + omp_in%x. - // - // These need to be temporary (in their own scope). If they are created - // as variables in the outer scope, if there's more than one type in the - // typelist, duplicate symbols will be reported. - const parser::CharBlock ompVarNames[]{ - {"omp_in", 6}, {"omp_out", 7}, {"omp_priv", 8}, {"omp_orig", 8}}; - - for (auto &t : typeList.v) { - PushScope(Scope::Kind::OtherConstruct, nullptr); - BeginDeclTypeSpec(); - // We need to walk t.u because Walk(t) does it's own BeginDeclTypeSpec. - Walk(t.u); + reductionDetails->AddDecl(declaratives_.back()); - // Only process types we can find. There will be an error later on when - // a type isn't found. - if (const DeclTypeSpec *typeSpec{GetDeclTypeSpec()}) { - reductionDetails->AddType(*typeSpec); + // Do not walk OmpTypeNameList. The types on the list will be visited + // during procesing of OmpCombinerExpression. + Walk(std::get<std::optional<parser::OmpCombinerExpression>>(spec.t)); + Walk(clauses); - for (auto &nm : ompVarNames) { - ObjectEntityDetails details{}; - details.set_type(*typeSpec); - MakeSymbol(nm, Attrs{}, std::move(details)); - } + for (auto &type : std::get<parser::OmpTypeNameList>(spec.t).v) { + // The declTypeSpec can be null if there is some semantic error. + if (type.declTypeSpec) { + reductionDetails->AddType(*type.declTypeSpec); } - EndDeclTypeSpec(); - Walk(std::get<std::optional<parser::OmpCombinerExpression>>(spec.t)); - Walk(clauses); - PopScope(); } - reductionDetails->AddDecl(declaratives_.back()); - if (!symbol) { symbol = &MakeSymbol(mangledName, Attrs{}, std::move(*reductionDetails)); } @@ -9456,13 +9435,18 @@ bool ResolveNamesVisitor::SetProcFlag( SayWithDecl(name, symbol, "Implicit declaration of function '%s' has a different result type than in previous declaration"_err_en_US); return false; - } else if (symbol.has<ProcEntityDetails>()) { - symbol.set(flag); // in case it hasn't been set yet - if (flag == Symbol::Flag::Function) { - ApplyImplicitRules(symbol); - } - if (symbol.attrs().test(Attr::INTRINSIC)) { - AcquireIntrinsicProcedureFlags(symbol); + } else if (const auto *proc{symbol.detailsIf<ProcEntityDetails>()}) { + if (IsPointer(symbol) && !proc->type() && !proc->procInterface()) { + // PROCEDURE(), POINTER -- errors will be emitted later about a lack + // of known characteristics if used as a function + } else { + symbol.set(flag); // in case it hasn't been set yet + if (flag == Symbol::Flag::Function) { + ApplyImplicitRules(symbol); + } + if (symbol.attrs().test(Attr::INTRINSIC)) { + AcquireIntrinsicProcedureFlags(symbol); + } } } else if (symbol.GetType() && flag == Symbol::Flag::Subroutine) { SayWithDecl( @@ -10130,6 +10114,9 @@ void ResolveNamesVisitor::Post(const parser::CompilerDirective &x) { case 'c': set.set(common::IgnoreTKR::Contiguous); break; + case 'p': + set.set(common::IgnoreTKR::Pointer); + break; case 'a': set = common::ignoreTKRAll; break; diff --git a/flang/lib/Support/Fortran.cpp b/flang/lib/Support/Fortran.cpp index 3a8ebbb..05d6e0e 100644 --- a/flang/lib/Support/Fortran.cpp +++ b/flang/lib/Support/Fortran.cpp @@ -95,6 +95,9 @@ std::string AsFortran(IgnoreTKRSet tkr) { if (tkr.test(IgnoreTKR::Contiguous)) { result += 'C'; } + if (tkr.test(IgnoreTKR::Pointer)) { + result += 'P'; + } return result; } |