diff options
Diffstat (limited to 'flang')
44 files changed, 1848 insertions, 290 deletions
diff --git a/flang/docs/Directives.md b/flang/docs/Directives.md index 3ebb08c..2f16a8d 100644 --- a/flang/docs/Directives.md +++ b/flang/docs/Directives.md @@ -1,9 +1,9 @@ -<!--===- docs/Directives.md - +<!--===- docs/Directives.md + 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 - + --> # Compiler directives supported by Flang @@ -12,16 +12,18 @@ A list of non-standard directives supported by Flang * `!dir$ fixed` and `!dir$ free` select Fortran source forms. Their effect persists to the end of the current source file. -* `!dir$ ignore_tkr [[(TKRDMAC)] dummy-arg-name]...` in an interface definition +* `!dir$ ignore_tkr [[(TKRDMACP)] dummy-arg-name]...` in an interface definition disables some semantic checks at call sites for the actual arguments that - correspond to some named dummy arguments (or all of them, by default). - The directive allow actual arguments that would otherwise be diagnosed - as incompatible in type (T), kind (K), rank (R), CUDA device (D), or - managed (M) status. The letter (A) is a shorthand for all of these, - and is the default when no letters appear. The letter (C) checks for - contiguity for example allowing an element of an assumed-shape array to be - passed as a dummy argument. For example, if one wanted to call a "set all - bytes to zero" utility that could be applied to arrays of any type or rank: + correspond to some named dummy arguments (or all of them, by default). The + directive allow actual arguments that would otherwise be diagnosed as + incompatible in type (T), kind (K), rank (R), CUDA device (D), or managed (M) + status. The letter (A) is a shorthand for (TKRDM), and is the default when no + letters appear. The letter (C) checks for contiguity, for example allowing an + element of an assumed-shape array to be passed as a dummy argument. The + letter (P) ignores pointer and allocatable matching, so that one can pass an + allocatable array to routine with pointer array argument and vice versa. For + example, if one wanted to call a "set all bytes to zero" utility that could + be applied to arrays of any type or rank: ``` interface subroutine clear(arr,bytes) @@ -46,27 +48,27 @@ A list of non-standard directives supported by Flang unroll the loop. Some compilers accept an optional `=` before the `n` when `n` is present in the directive. Flang does not. * `!dir$ unroll_and_jam [N]` control how many times a loop should be unrolled and - jammed. It must be placed immediately before a loop that follows. `N` is an optional - integer that specifying the unrolling factor. When `N` is `0` or `1`, the loop + jammed. It must be placed immediately before a loop that follows. `N` is an optional + integer that specifying the unrolling factor. When `N` is `0` or `1`, the loop should not be unrolled at all. If `N` is omitted the optimizer will selects the number of times to unroll the loop. * `!dir$ novector` disabling vectorization on the following loop. * `!dir$ nounroll` disabling unrolling on the following loop. * `!dir$ nounroll_and_jam` disabling unrolling and jamming on the following loop. -* `!dir$ inline` instructs the compiler to attempt to inline the called routines if the - directive is specified before a call statement or all call statements within the loop - body if specified before a DO LOOP or all function references if specified before an +* `!dir$ inline` instructs the compiler to attempt to inline the called routines if the + directive is specified before a call statement or all call statements within the loop + body if specified before a DO LOOP or all function references if specified before an assignment statement. -* `!dir$ forceinline` works in the same way as the `inline` directive, but it forces +* `!dir$ forceinline` works in the same way as the `inline` directive, but it forces inlining by the compiler on a function call statement. -* `!dir$ noinline` works in the same way as the `inline` directive, but prevents +* `!dir$ noinline` works in the same way as the `inline` directive, but prevents any attempt of inlining by the compiler on a function call statement. # Directive Details ## Introduction -Directives are commonly used in Fortran programs to specify additional actions -to be performed by the compiler. The directives are always specified with the +Directives are commonly used in Fortran programs to specify additional actions +to be performed by the compiler. The directives are always specified with the `!dir$` or `cdir$` prefix. ## Loop Directives @@ -97,7 +99,7 @@ check that that construct matches the expected construct for the directive. Skipping other intermediate directives allows multiple directives to appear on the same construct. -## Lowering +## Lowering Evaluation is extended with a new field called dirs for representing directives associated with that Evaluation. When lowering loop directives, the associated Do Loop's evaluation is found and the directive is added to it. This information @@ -109,7 +111,7 @@ about the loop. For example, the `llvm.loop.vectorize.enable` metadata informs the optimizer that a loop can be vectorized without considering its cost-model. This attribute is added to the loop condition branch. -### Representation in MLIR +### Representation in MLIR The MLIR LLVM dialect models this by an attribute called LoopAnnotation Attribute. The attribute can be added to the latch of the loop in the cf dialect and is then carried through lowering to the LLVM dialect. diff --git a/flang/include/flang/Optimizer/Builder/IntrinsicCall.h b/flang/include/flang/Optimizer/Builder/IntrinsicCall.h index c3cd119b..3407dd0 100644 --- a/flang/include/flang/Optimizer/Builder/IntrinsicCall.h +++ b/flang/include/flang/Optimizer/Builder/IntrinsicCall.h @@ -211,6 +211,8 @@ struct IntrinsicLibrary { mlir::Value genBarrierArrive(mlir::Type, llvm::ArrayRef<mlir::Value>); mlir::Value genBarrierArriveCnt(mlir::Type, llvm::ArrayRef<mlir::Value>); void genBarrierInit(llvm::ArrayRef<fir::ExtendedValue>); + mlir::Value genBarrierTryWait(mlir::Type, llvm::ArrayRef<mlir::Value>); + mlir::Value genBarrierTryWaitSleep(mlir::Type, llvm::ArrayRef<mlir::Value>); fir::ExtendedValue genBesselJn(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>); fir::ExtendedValue genBesselYn(mlir::Type, @@ -459,7 +461,21 @@ struct IntrinsicLibrary { mlir::Value genTime(mlir::Type, llvm::ArrayRef<mlir::Value>); void genTMABulkCommitGroup(llvm::ArrayRef<fir::ExtendedValue>); void genTMABulkG2S(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkLoadC4(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkLoadC8(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkLoadI4(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkLoadI8(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkLoadR2(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkLoadR4(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkLoadR8(llvm::ArrayRef<fir::ExtendedValue>); void genTMABulkS2G(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkStoreI4(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkStoreI8(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkStoreR2(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkStoreR4(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkStoreR8(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkStoreC4(llvm::ArrayRef<fir::ExtendedValue>); + void genTMABulkStoreC8(llvm::ArrayRef<fir::ExtendedValue>); void genTMABulkWaitGroup(llvm::ArrayRef<fir::ExtendedValue>); mlir::Value genTrailz(mlir::Type, llvm::ArrayRef<mlir::Value>); fir::ExtendedValue genTransfer(mlir::Type, diff --git a/flang/include/flang/Parser/dump-parse-tree.h b/flang/include/flang/Parser/dump-parse-tree.h index 553cbd5..bb97069 100644 --- a/flang/include/flang/Parser/dump-parse-tree.h +++ b/flang/include/flang/Parser/dump-parse-tree.h @@ -599,7 +599,7 @@ public: NODE(parser, OmpInitClause) NODE(OmpInitClause, Modifier) NODE(parser, OmpInitializerClause) - NODE(parser, OmpInitializerProc) + NODE(parser, OmpInitializerExpression) NODE(parser, OmpInReductionClause) NODE(OmpInReductionClause, Modifier) NODE(parser, OmpInteropPreference) @@ -677,6 +677,10 @@ public: NODE_ENUM(OmpSeverityClause, Severity) NODE(parser, OmpStepComplexModifier) NODE(parser, OmpStepSimpleModifier) + NODE(parser, OmpStylizedDeclaration) + NODE(parser, OmpStylizedExpression) + NODE(parser, OmpStylizedInstance) + NODE(OmpStylizedInstance, Instance) NODE(parser, OmpTaskDependenceType) NODE_ENUM(OmpTaskDependenceType, Value) NODE(parser, OmpTaskReductionClause) diff --git a/flang/include/flang/Parser/openmp-utils.h b/flang/include/flang/Parser/openmp-utils.h index f761332..49db091 100644 --- a/flang/include/flang/Parser/openmp-utils.h +++ b/flang/include/flang/Parser/openmp-utils.h @@ -25,6 +25,13 @@ namespace Fortran::parser::omp { +template <typename T> constexpr auto addr_if(std::optional<T> &x) { + return x ? &*x : nullptr; +} +template <typename T> constexpr auto addr_if(const std::optional<T> &x) { + return x ? &*x : nullptr; +} + namespace detail { using D = llvm::omp::Directive; @@ -133,9 +140,24 @@ template <typename T> OmpDirectiveName GetOmpDirectiveName(const T &x) { } const OmpObjectList *GetOmpObjectList(const OmpClause &clause); + +template <typename T> +const T *GetFirstArgument(const OmpDirectiveSpecification &spec) { + for (const OmpArgument &arg : spec.Arguments().v) { + if (auto *t{std::get_if<T>(&arg.u)}) { + return t; + } + } + return nullptr; +} + const BlockConstruct *GetFortranBlockConstruct( const ExecutionPartConstruct &epc); +const OmpCombinerExpression *GetCombinerExpr( + const OmpReductionSpecifier &rspec); +const OmpInitializerExpression *GetInitializerExpr(const OmpClause &init); + } // namespace Fortran::parser::omp #endif // FORTRAN_PARSER_OPENMP_UTILS_H diff --git a/flang/include/flang/Parser/parse-tree.h b/flang/include/flang/Parser/parse-tree.h index 2cf6fae..c3a8c2e 100644 --- a/flang/include/flang/Parser/parse-tree.h +++ b/flang/include/flang/Parser/parse-tree.h @@ -24,7 +24,9 @@ #include "provenance.h" #include "flang/Common/idioms.h" #include "flang/Common/indirection.h" +#include "flang/Common/reference.h" #include "flang/Support/Fortran.h" +#include "llvm/ADT/ArrayRef.h" #include "llvm/Frontend/OpenACC/ACC.h.inc" #include "llvm/Frontend/OpenMP/OMP.h" #include "llvm/Frontend/OpenMP/OMPConstants.h" @@ -3510,6 +3512,8 @@ struct OmpDirectiveName { // type-name list item struct OmpTypeName { + CharBlock source; + mutable const semantics::DeclTypeSpec *declTypeSpec{nullptr}; UNION_CLASS_BOILERPLATE(OmpTypeName); std::variant<TypeSpec, DeclarationTypeSpec> u; }; @@ -3538,6 +3542,39 @@ struct OmpObjectList { WRAPPER_CLASS_BOILERPLATE(OmpObjectList, std::list<OmpObject>); }; +struct OmpStylizedDeclaration { + COPY_AND_ASSIGN_BOILERPLATE(OmpStylizedDeclaration); + // Since "Reference" isn't handled by parse-tree-visitor, add EmptyTrait, + // and visit the members by hand when needed. + using EmptyTrait = std::true_type; + common::Reference<const OmpTypeName> type; + EntityDecl var; +}; + +struct OmpStylizedInstance { + struct Instance { + UNION_CLASS_BOILERPLATE(Instance); + std::variant<AssignmentStmt, CallStmt, common::Indirection<Expr>> u; + }; + TUPLE_CLASS_BOILERPLATE(OmpStylizedInstance); + std::tuple<std::list<OmpStylizedDeclaration>, Instance> t; +}; + +class ParseState; + +// Ref: [5.2:76], [6.0:185] +// +struct OmpStylizedExpression { + CharBlock source; + // Pointer to a temporary copy of the ParseState that is used to create + // additional parse subtrees for the stylized expression. This is only + // used internally during parsing and conveys no information to the + // consumers of the AST. + const ParseState *state{nullptr}; + WRAPPER_CLASS_BOILERPLATE( + OmpStylizedExpression, std::list<OmpStylizedInstance>); +}; + // Ref: [4.5:201-207], [5.0:293-299], [5.1:325-331], [5.2:124] // // reduction-identifier -> @@ -3555,9 +3592,22 @@ struct OmpReductionIdentifier { // combiner-expression -> // since 4.5 // assignment-statement | // function-reference -struct OmpCombinerExpression { - UNION_CLASS_BOILERPLATE(OmpCombinerExpression); - std::variant<AssignmentStmt, FunctionReference> u; +struct OmpCombinerExpression : public OmpStylizedExpression { + INHERITED_WRAPPER_CLASS_BOILERPLATE( + OmpCombinerExpression, OmpStylizedExpression); + static llvm::ArrayRef<CharBlock> Variables(); +}; + +// Ref: [4.5:222:7-8], [5.0:305:28-29], [5.1:337:20-21], [5.2:127:6-8], +// [6.0:242:3-5] +// +// initializer-expression -> // since 4.5 +// OMP_PRIV = expression | +// subroutine-name(argument-list) +struct OmpInitializerExpression : public OmpStylizedExpression { + INHERITED_WRAPPER_CLASS_BOILERPLATE( + OmpInitializerExpression, OmpStylizedExpression); + static llvm::ArrayRef<CharBlock> Variables(); }; inline namespace arguments { @@ -4558,16 +4608,9 @@ struct OmpInReductionClause { std::tuple<MODIFIERS(), OmpObjectList> t; }; -// declare-reduction -> DECLARE REDUCTION (reduction-identifier : type-list -// : combiner) [initializer-clause] -struct OmpInitializerProc { - TUPLE_CLASS_BOILERPLATE(OmpInitializerProc); - std::tuple<ProcedureDesignator, std::list<ActualArgSpec>> t; -}; // Initialization for declare reduction construct struct OmpInitializerClause { - UNION_CLASS_BOILERPLATE(OmpInitializerClause); - std::variant<OmpInitializerProc, AssignmentStmt> u; + WRAPPER_CLASS_BOILERPLATE(OmpInitializerClause, OmpInitializerExpression); }; // Ref: [4.5:199-201], [5.0:288-290], [5.1:321-322], [5.2:115-117] diff --git a/flang/include/flang/Semantics/symbol.h b/flang/include/flang/Semantics/symbol.h index 04a0639..cb27d544 100644 --- a/flang/include/flang/Semantics/symbol.h +++ b/flang/include/flang/Semantics/symbol.h @@ -830,6 +830,8 @@ public: OmpUseDevicePtr, OmpUseDeviceAddr, OmpIsDevicePtr, OmpHasDeviceAddr, // OpenMP data-copying attribute OmpCopyIn, OmpCopyPrivate, + // OpenMP special variables + OmpInVar, OmpOrigVar, OmpOutVar, OmpPrivVar, // OpenMP miscellaneous flags OmpCommonBlock, OmpReduction, OmpInReduction, OmpAligned, OmpNontemporal, OmpAllocate, OmpDeclarativeAllocateDirective, diff --git a/flang/include/flang/Support/Fortran.h b/flang/include/flang/Support/Fortran.h index ea0344e..cf39781 100644 --- a/flang/include/flang/Support/Fortran.h +++ b/flang/include/flang/Support/Fortran.h @@ -86,8 +86,9 @@ ENUM_CLASS(IgnoreTKR, Rank, // R - don't check ranks Device, // D - don't check host/device residence Managed, // M - don't check managed storage - Contiguous) // C - don't check for storage sequence association with a + Contiguous, // C - don't check for storage sequence association with a // potentially non-contiguous object + Pointer) // P - ignore pointer and allocatable matching using IgnoreTKRSet = EnumSet<IgnoreTKR, 8>; // IGNORE_TKR(A) = IGNORE_TKR(TKRDM) static constexpr IgnoreTKRSet ignoreTKRAll{IgnoreTKR::Type, IgnoreTKR::Kind, diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp index 39bac81..ca3e1cd 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, {{}}, @@ -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/Dialect/FIROps.cpp b/flang/lib/Optimizer/Dialect/FIROps.cpp index d0164f3..4f97aca 100644 --- a/flang/lib/Optimizer/Dialect/FIROps.cpp +++ b/flang/lib/Optimizer/Dialect/FIROps.cpp @@ -4484,7 +4484,7 @@ void fir::IfOp::getSuccessorRegions( llvm::SmallVectorImpl<mlir::RegionSuccessor> ®ions) { // The `then` and the `else` region branch back to the parent operation. if (!point.isParent()) { - regions.push_back(mlir::RegionSuccessor(getResults())); + regions.push_back(mlir::RegionSuccessor(getOperation(), getResults())); return; } @@ -4494,7 +4494,8 @@ void fir::IfOp::getSuccessorRegions( // Don't consider the else region if it is empty. mlir::Region *elseRegion = &this->getElseRegion(); if (elseRegion->empty()) - regions.push_back(mlir::RegionSuccessor()); + regions.push_back( + mlir::RegionSuccessor(getOperation(), getOperation()->getResults())); else regions.push_back(mlir::RegionSuccessor(elseRegion)); } @@ -4513,7 +4514,7 @@ void fir::IfOp::getEntrySuccessorRegions( if (!getElseRegion().empty()) regions.emplace_back(&getElseRegion()); else - regions.emplace_back(getResults()); + regions.emplace_back(getOperation(), getOperation()->getResults()); } } 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/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-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-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/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..f88af5f 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)); } @@ -10130,6 +10109,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; } diff --git a/flang/module/cudadevice.f90 b/flang/module/cudadevice.f90 index 5182950..59af58d 100644 --- a/flang/module/cudadevice.f90 +++ b/flang/module/cudadevice.f90 @@ -1998,6 +1998,18 @@ implicit none ! TMA Operations + interface barrier_arrive + attributes(device) function barrier_arrive(barrier) result(token) + integer(8), shared :: barrier + integer(8) :: token + end function + attributes(device) function barrier_arrive_cnt(barrier, count) result(token) + integer(8), shared :: barrier + integer(4), value :: count + integer(8) :: token + end function + end interface + interface attributes(device) subroutine barrier_init(barrier, count) integer(8), shared :: barrier @@ -2005,15 +2017,18 @@ implicit none end subroutine end interface - interface barrier_arrive - attributes(device) function barrier_arrive(barrier) result(token) + interface + attributes(device) integer function barrier_try_wait(barrier, token) integer(8), shared :: barrier - integer(8) :: token + integer(8), value :: token end function - attributes(device) function barrier_arrive_cnt(barrier, count) result(token) + end interface + + interface + attributes(device) integer function barrier_try_wait_sleep(barrier, token, ns) integer(8), shared :: barrier - integer(4), value :: count - integer(8) :: token + integer(8), value :: token + integer(4), value :: ns end function end interface @@ -2032,7 +2047,13 @@ implicit none end subroutine end interface + ! -------------------- + ! Bulk load functions + ! -------------------- + ! Generic load, count is in bytes + ! ------------------------------- + interface attributes(device) subroutine tma_bulk_g2s(barrier, src, dst, nbytes) !dir$ ignore_tkr src, dst @@ -2043,6 +2064,74 @@ implicit none end subroutine end interface + ! Load specific types, count is in elements + ! ----------------------------------------- + + interface tma_bulk_load + attributes(device) subroutine tma_bulk_ldc4(barrier, src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: barrier + complex(4), device :: src(*) + complex(4), shared :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_ldc8(barrier, src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: barrier + complex(8), device :: src(*) + complex(8), shared :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_ldi4(barrier, src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: barrier + integer(4), device :: src(*) + integer(4), shared :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_ldi8(barrier, src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: barrier + integer(8), device :: src(*) + integer(8), shared :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_ldr2(barrier, src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: barrier + real(2), device :: src(*) + real(2), shared :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_ldr4(barrier, src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: barrier + real(4), device :: src(*) + real(4), shared :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_ldr8(barrier, src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: barrier + real(8), device :: src(*) + real(8), shared :: dst(*) + integer(4), value :: nelems + end subroutine + end interface + + ! -------------------- + ! Bulk Store functions + ! -------------------- + + ! Generic store, count is in bytes + ! -------------------------------- + interface attributes(device) subroutine tma_bulk_s2g(src, dst, nbytes) !dir$ ignore_tkr src, dst @@ -2052,6 +2141,60 @@ implicit none end subroutine end interface + ! Load specific types, count is in elements + ! ----------------------------------------- + + interface tma_bulk_store + attributes(device) subroutine tma_bulk_store_c4(src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + complex(4), shared :: src(*) + complex(4), device :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_store_c8(src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + complex(8), shared :: src(*) + complex(8), device :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_store_i4(src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(4), shared :: src(*) + integer(4), device :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_store_i8(src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + integer(8), shared :: src(*) + integer(8), device :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_store_r2(src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + real(2), shared :: src(*) + real(2), device :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_store_r4(src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + real(4), shared :: src(*) + real(4), device :: dst(*) + integer(4), value :: nelems + end subroutine + + attributes(device) subroutine tma_bulk_store_r8(src, dst, nelems) + !dir$ ignore_tkr (r) src, (r) dst + real(8), shared :: src(*) + real(8), device :: dst(*) + integer(4), value :: nelems + end subroutine + end interface + contains attributes(device) subroutine syncthreads() diff --git a/flang/test/Driver/flang-f-opts.f90 b/flang/test/Driver/flang-f-opts.f90 index 77bb4d7..9ef0aba 100644 --- a/flang/test/Driver/flang-f-opts.f90 +++ b/flang/test/Driver/flang-f-opts.f90 @@ -1,5 +1,5 @@ -! Test for warnings generated when parsing driver options. You can use this file for relatively small tests and to avoid creating -! new test files. +! Test for errors and warnings generated when parsing driver options. You can +! use this file for relatively small tests and to avoid creating new test files. ! RUN: %flang -### -S -O4 -ffp-contract=on %s 2>&1 | FileCheck %s @@ -26,3 +26,20 @@ ! RUN: | FileCheck %s -check-prefix=WARN-BUILTIN-MULTIPLE ! WARN-BUILTIN-MULTIPLE: warning: '-fbuiltin' is not valid for Fortran ! WARN-BUILTIN-MULTIPLE: warning: '-fno-builtin' is not valid for Fortran + +! When emitting an error with a suggestion, ensure that the diagnostic message +! uses '-Xflang' instead of '-Xclang'. This is typically emitted when an option +! that is available for `flang -fc1` is passed to `flang`. We use -complex-range +! since it is only available for fc1. If this option is ever exposed to `flang`, +! a different option will have to be used in the test below. +! +! RUN: not %flang -### -complex-range=full %s 2>&1 \ +! RUN: | FileCheck %s -check-prefix UNKNOWN-SUGGEST +! +! UNKNOWN-SUGGEST: error: unknown argument '-complex-range=full'; +! UNKNOWN-SUGGEST-SAME: did you mean '-Xflang -complex-range=full' +! +! RUN: not %flang -### -not-an-option %s 2>&1 \ +! RUN: | FileCheck %s -check-prefix UNKNOWN-NO-SUGGEST +! +! UNKNOWN-NO-SUGGEST: error: unknown argument: '-not-an-option'{{$}} diff --git a/flang/test/Driver/linker-options.f90 b/flang/test/Driver/linker-options.f90 new file mode 100644 index 0000000..07f967b --- /dev/null +++ b/flang/test/Driver/linker-options.f90 @@ -0,0 +1,106 @@ +! Make sure that `-l` is "visible" to Flang's driver +! RUN: %flang -lpgmath -### %s + +! Make sure that `-Wl` is "visible" to Flang's driver +! RUN: %flang -Wl,abs -### %s + +! Make sure that `-fuse-ld' is "visible" to Flang's driver +! RUN: %flang -fuse-ld= -### %s + +! Make sure that `-L' is "visible" to Flang's driver +! RUN: %flang -L/ -### %s + +! ------------------------------------------------------------------------------ +! Check that '-pie' and '-no-pie' are "visible" to Flang's driver. Check that +! the correct option is added to the link line. +! +! Last match "wins" +! RUN: %flang -target x86_64-pc-linux-gnu -pie -no-pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target x86_64-pc-linux-gnu -no-pie -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! RUN: %flang -target x86_64-pc-linux-gnu -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! RUN: %flang -target x86_64-pc-linux-gnu -no-pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! +! Ensure that "-pie" is passed to the linker. +! RUN: %flang -target i386-unknown-freebsd -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! RUN: %flang -target aarch64-pc-linux-gnu -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! +! On Musl Linux, PIE is enabled by default, but can be disabled. +! RUN: %flang -target x86_64-linux-musl -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! RUN: %flang -target i686-linux-musl -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! RUN: %flang -target armv6-linux-musleabihf %s -### 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! RUN: %flang -target armv7-linux-musleabihf %s -### 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! RUN: %flang --target=x86_64-linux-musl -no-pie -### 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! +! On OpenBSD, -pie is not passed to the linker, but can be forced. +! RUN: %flang -target amd64-pc-openbsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target i386-pc-openbsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target aarch64-unknown-openbsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target arm-unknown-openbsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target powerpc-unknown-openbsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target sparc64-unknown-openbsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target i386-pc-openbsd -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! +! On FreeBSD, -pie is not passed to the linker, but can be forced. +! RUN: %flang -target amd64-pc-freebsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target i386-pc-freebsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target aarch64-unknown-freebsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target arm-unknown-freebsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target powerpc-unknown-freebsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target sparc64-unknown-freebsd -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=NO-PIE +! RUN: %flang -target i386-pc-freebsd -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefix=PIE +! +! On AIX, -pie is never passed to the linker. +! RUN: %flang -target powerpc64-unknown-aix -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE +! RUN: %flang -target powerpc64-unknown-aix -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE,UNUSED +! RUN: %flang -target powerpc64-unknown-aix -no-pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE,UNUSED +! +! On MinGW and Windows, -pie may be specified, but it is ignored. +! RUN: %flang -target aarch64-pc-windows-gnu -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE +! RUN: %flang -target x86_64-pc-windows-gnu -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE,UNUSED +! RUN: %flang -target i686-pc-windows-gnu -no-pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE,UNUSED +! RUN: %flang -target aarch64-windows-msvc -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE +! RUN: %flang -target aarch64-windows-msvc -pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE,UNUSED +! RUN: %flang -target aarch64-windows-msvc -no-pie -### %s 2>&1 \ +! RUN: | FileCheck %s --check-prefixes=NO-PIE,UNUSED +! +! PIE: "-pie" +! NO-PIE-NOT: "-pie" +! UNUSED: warning: argument unused during compilation: '{{(-no)?}}-pie' +! ------------------------------------------------------------------------------ + +program hello + write(*,*), "Hello world!" +end program hello diff --git a/flang/test/Driver/misc-flags.f90 b/flang/test/Driver/misc-flags.f90 deleted file mode 100644 index 61d763c..0000000 --- a/flang/test/Driver/misc-flags.f90 +++ /dev/null @@ -1,15 +0,0 @@ -! Make sure that `-l` is "visible" to Flang's driver -! RUN: %flang -lpgmath -### %s - -! Make sure that `-Wl` is "visible" to Flang's driver -! RUN: %flang -Wl,abs -### %s - -! Make sure that `-fuse-ld' is "visible" to Flang's driver -! RUN: %flang -fuse-ld= -### %s - -! Make sure that `-L' is "visible" to Flang's driver -! RUN: %flang -L/ -### %s - -program hello - write(*,*), "Hello world!" -end program hello diff --git a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir b/flang/test/Fir/CUDA/cuda-target-rewrite.mlir index 48fee10..5562e00 100644 --- a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir +++ b/flang/test/Fir/CUDA/cuda-target-rewrite.mlir @@ -108,3 +108,23 @@ module attributes {gpu.container_module, fir.defaultkind = "a1c4d8i4l4r4", fir.k } } +// ----- + +module attributes {gpu.container_module, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.target_triple = "x86_64-unknown-linux-gnu"} { + gpu.module @testmod { + gpu.func @_QPtest(%arg0: complex<f32>) -> () kernel { + gpu.return + } + } + func.func @main(%arg0: complex<f32>) { + %0 = llvm.mlir.constant(0 : i64) : i64 + %1 = llvm.mlir.constant(0 : i32) : i32 + %2 = fir.alloca i64 + %3 = cuf.stream_cast %2 : !fir.ref<i64> + %4 = gpu.launch_func async [%3] @testmod::@_QPtest blocks in (%0, %0, %0) threads in (%0, %0, %0) : i64 dynamic_shared_memory_size %1 args(%arg0 : complex<f32>) {cuf.proc_attr = #cuf.cuda_proc<global>} + return + } +} + +// CHECK-LABEL: func.func @main +// CHECK: %{{.*}} = gpu.launch_func async [%{{.*}}] @testmod::@_QPtest blocks in (%{{.*}}, %{{.*}}, %{{.*}}) threads in (%{{.*}}, %{{.*}}, %{{.*}}) : i64 dynamic_shared_memory_size %{{.*}} args(%{{.*}} : !fir.vector<2:f32>) {cuf.proc_attr = #cuf.cuda_proc<global>} diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf index 5c4c3c6..e5d3c43 100644 --- a/flang/test/Lower/CUDA/cuda-device-proc.cuf +++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf @@ -468,7 +468,18 @@ attributes(global) subroutine test_bulk_g2s(a) end subroutine ! CHECK-LABEL: func.func @_QPtest_bulk_g2s -! CHECK: nvvm.cp.async.bulk.shared.cluster.global %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : <7>, <1> +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %4 {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEbarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[DST:.*]]:2 = hlfir.declare %16(%17) {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEtmpa"} : (!fir.ref<!fir.array<1024xf64>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf64>>, !fir.ref<!fir.array<1024xf64>>) +! CHECK: %[[COUNT:.*]]:2 = hlfir.declare %19 {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_bulk_g2sEtx_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[SRC:.*]] = hlfir.designate %{{.*}} (%{{.*}}) : (!fir.box<!fir.array<?xf64>>, i64) -> !fir.ref<f64> +! CHECK: %[[COUNT_LOAD:.*]] = fir.load %20#0 : !fir.ref<i32> +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: %[[BARRIER_3:.*]] = llvm.addrspacecast %[[BARRIER_PTR]] : !llvm.ptr to !llvm.ptr<3> +! CHECK: %[[DST_PTR:.*]] = fir.convert %[[DST]]#0 : (!fir.ref<!fir.array<1024xf64>>) -> !llvm.ptr +! CHECK: %[[DST_7:.*]] = llvm.addrspacecast %[[DST_PTR]] : !llvm.ptr to !llvm.ptr<7> +! CHECK: %[[SRC_PTR:.*]] = fir.convert %[[SRC]] : (!fir.ref<f64>) -> !llvm.ptr +! CHECK: %[[SRC_3:.*]] = llvm.addrspacecast %[[SRC_PTR]] : !llvm.ptr to !llvm.ptr<1> +! CHECK: nvvm.cp.async.bulk.shared.cluster.global %[[DST_7]], %[[SRC_3]], %[[BARRIER_3]], %[[COUNT_LOAD]] : <7>, <1> attributes(global) subroutine test_bulk_s2g(a) real(8), device :: a(*) @@ -479,6 +490,8 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_bulk_s2g ! CHECL: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(device) subroutine testAtomicCasLoop(aa, n) integer :: a @@ -492,3 +505,250 @@ end subroutine ! CHECK: %[[CASTED_CMP_XCHG_EV:.*]] = fir.convert %[[CMP_XCHG_EV]] : (i1) -> i32 ! CHECK: %{{.*}} = arith.constant 1 : i32 ! CHECK: %19 = arith.cmpi eq, %[[CASTED_CMP_XCHG_EV]], %{{.*}} : i32 + +attributes(global) subroutine test_barrier_try_wait() + integer :: istat + integer(8), shared :: barrier1 + integer(8) :: token + istat = barrier_try_wait(barrier1, token) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_barrier_try_wait() +! CHECK: scf.while +! CHECK: %{{.*}} = nvvm.inline_ptx ".reg .pred p; mbarrier.try_wait.shared.b64 p, [%{{.*}}], %{{.*}}, %{{.*}}; selp.b32 %{{.*}}, 1, 0, p;" ro(%{{.*}}, %{{.*}}, %c1000000{{.*}} : !llvm.ptr, i64, i32) -> i32 + +attributes(global) subroutine test_barrier_try_wait_sleep() + integer :: istat + integer(8), shared :: barrier1 + integer(8) :: token + integer(4) :: sleep_time + istat = barrier_try_wait_sleep(barrier1, token, sleep_time) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_barrier_try_wait_sleep() +! CHECK: %{{.*}} = nvvm.inline_ptx ".reg .pred p; mbarrier.try_wait.shared.b64 p, [%{{.*}}], %{{.*}}, %{{.*}}; selp.b32 %0, 1, 0, p;" ro(%{{.*}}, %{{.*}}, %{{.*}} : !llvm.ptr, i64, i32) -> i32 + +attributes(global) subroutine test_tma_bulk_load_c4(a, n) + integer(8), shared :: barrier1 + integer, value :: n + complex(4), device :: r8(n) + complex(4), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_c4 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_c4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_c4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_c8(a, n) + integer(8), shared :: barrier1 + integer, value :: n + complex(8), device :: r8(n) + complex(8), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_c8 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_c8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_c8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 16 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_i4(a, n) + integer(8), shared :: barrier1 + integer, value :: n + integer(4), device :: r8(n) + integer(4), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_i4 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_i4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_i4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_i8(a, n) + integer(8), shared :: barrier1 + integer, value :: n + integer(8), device :: r8(n) + integer(8), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_i8 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_i8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_i8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_r2(a, n) + integer(8), shared :: barrier1 + integer, value :: n + real(2), device :: r8(n) + real(2), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r2 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r2Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r2Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 2 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_r4(a, n) + integer(8), shared :: barrier1 + integer, value :: n + real(4), device :: r8(n) + real(4), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r4 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_r8(a, n) + integer(8), shared :: barrier1 + integer, value :: n + real(8), device :: r8(n) + real(8), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r8 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_store_c4(c, n) + integer, value :: n + complex(4), device :: c(n) + complex(4), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c4 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_c8(c, n) + integer, value :: n + complex(8), device :: c(n) + complex(8), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c8 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_i4(c, n) + integer, value :: n + integer(4), device :: c(n) + integer(4), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i4 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_i8(c, n) + integer, value :: n + integer(8), device :: c(n) + integer(8), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i8 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + + +attributes(global) subroutine test_tma_bulk_store_r2(c, n) + integer, value :: n + real(2), device :: c(n) + real(2), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r2 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_r4(c, n) + integer, value :: n + real(4), device :: c(n) + real(4), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r4 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_r8(c, n) + integer, value :: n + real(8), device :: c(n) + real(8), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r8 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 diff --git a/flang/test/Lower/OpenMP/atomic-read-complex.f90 b/flang/test/Lower/OpenMP/atomic-read-complex.f90 new file mode 100644 index 0000000..2f51f03 --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-read-complex.f90 @@ -0,0 +1,34 @@ +! Test lowering of atomic read to LLVM IR for complex types. +! This is a regression test for issue #165184. + +! RUN: %flang_fc1 -emit-llvm -fopenmp -o - %s | FileCheck %s + +! Test that atomic read operations with complex types emit the correct +! size parameter to __atomic_load: +! - complex(4) (8 bytes total): should call __atomic_load(i64 8, ...) +! - complex(8) (16 bytes total): should call __atomic_load(i64 16, ...) + +program atomic_read_complex + implicit none + + ! Test complex(4) - single precision (8 bytes) + complex(4) :: c41, c42 + ! Test complex(8) - double precision (16 bytes) + complex(8) :: c81, c82 + + c42 = (1.0_4, 1.0_4) + c82 = (1.0_8, 1.0_8) + + ! CHECK-LABEL: define {{.*}} @_QQmain + + ! Single precision complex: 8 bytes + ! CHECK: call void @__atomic_load(i64 8, ptr {{.*}}, ptr {{.*}}, i32 {{.*}}) +!$omp atomic read + c41 = c42 + + ! Double precision complex: 16 bytes (this was broken before the fix) + ! CHECK: call void @__atomic_load(i64 16, ptr {{.*}}, ptr {{.*}}, i32 {{.*}}) +!$omp atomic read + c81 = c82 + +end program atomic_read_complex diff --git a/flang/test/Lower/OpenMP/atomic-write-complex.f90 b/flang/test/Lower/OpenMP/atomic-write-complex.f90 new file mode 100644 index 0000000..48cfe26 --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-write-complex.f90 @@ -0,0 +1,34 @@ +! Test lowering of atomic write to LLVM IR for complex types. +! This is a regression test for issue #165184. + +! RUN: %flang_fc1 -emit-llvm -fopenmp -o - %s | FileCheck %s + +! Test that atomic write operations with complex types emit the correct +! size parameter to __atomic_store: +! - complex(4) (8 bytes total): should call __atomic_store(i64 8, ...) +! - complex(8) (16 bytes total): should call __atomic_store(i64 16, ...) + +program atomic_write_complex + implicit none + + ! Test complex(4) - single precision (8 bytes) + complex(4) :: c41, c42 + ! Test complex(8) - double precision (16 bytes) + complex(8) :: c81, c82 + + c42 = (1.0_4, 1.0_4) + c82 = (1.0_8, 1.0_8) + + ! CHECK-LABEL: define {{.*}} @_QQmain + + ! Single precision complex: 8 bytes + ! CHECK: call void @__atomic_store(i64 8, ptr {{.*}}, ptr {{.*}}, i32 {{.*}}) +!$omp atomic write + c41 = c42 + + ! Double precision complex: 16 bytes (this was broken before the fix) + ! CHECK: call void @__atomic_store(i64 16, ptr {{.*}}, ptr {{.*}}, i32 {{.*}}) +!$omp atomic write + c81 = c82 + +end program atomic_write_complex diff --git a/flang/test/Parser/OpenMP/declare-reduction-multi.f90 b/flang/test/Parser/OpenMP/declare-reduction-multi.f90 index a682958..8856661 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-multi.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-multi.f90 @@ -26,7 +26,8 @@ program omp_examples type(tt) :: values(n), sum, prod, big, small !$omp declare reduction(+:tt:omp_out%r = omp_out%r + omp_in%r) initializer(omp_priv%r = 0) -!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out%r = omp_out%r+omp_in%r) INITIALIZER(omp_priv%r = 0_4) +!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out%r = omp_out%r + omp_in%r) INITIALIZER(om& +!CHECK-NEXT: !$OMP&p_priv%r = 0) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -34,11 +35,39 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%r=omp_out%r+omp_in%r' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=0._4' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out%r=omp_out%r+omp_in%r' +!PARSE-TREE: | | | | Variable = 'omp_out%r' +!PARSE-TREE: | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Name = 'r' +!PARSE-TREE: | | | | Expr = 'omp_out%r+omp_in%r' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | | | | | | Expr = 'omp_in%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=0._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '0_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$omp declare reduction(*:tt:omp_out%r = omp_out%r * omp_in%r) initializer(omp_priv%r = 1) -!CHECK-NEXT: !$OMP DECLARE REDUCTION(*:tt: omp_out%r = omp_out%r*omp_in%r) INITIALIZER(omp_priv%r = 1_4) +!CHECK-NEXT: !$OMP DECLARE REDUCTION(*:tt: omp_out%r = omp_out%r * omp_in%r) INITIALIZER(om& +!CHECK-NEXT: !$OMP&p_priv%r = 1) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -46,11 +75,39 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Multiply !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%r=omp_out%r*omp_in%r' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=1._4' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out%r=omp_out%r*omp_in%r' +!PARSE-TREE: | | | | Variable = 'omp_out%r' +!PARSE-TREE: | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Name = 'r' +!PARSE-TREE: | | | | Expr = 'omp_out%r*omp_in%r' +!PARSE-TREE: | | | | | Multiply +!PARSE-TREE: | | | | | | Expr = 'omp_out%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | | | | | | Expr = 'omp_in%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=1._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '1_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '1' +!PARSE-TREE: | Flags = None !$omp declare reduction(max:tt:omp_out = mymax(omp_out, omp_in)) initializer(omp_priv%r = 0) -!CHECK-NEXT: !$OMP DECLARE REDUCTION(max:tt: omp_out = mymax(omp_out,omp_in)) INITIALIZER(omp_priv%r = 0_4) +!CHECK-NEXT: !$OMP DECLARE REDUCTION(max:tt: omp_out = mymax(omp_out, omp_in)) INITIALIZER(& +!CHECK-NEXT: !$OMP&omp_priv%r = 0) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -58,11 +115,36 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'max' !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=mymax(omp_out,omp_in)' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=0._4' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=mymax(omp_out,omp_in)' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'mymax(omp_out,omp_in)' +!PARSE-TREE: | | | | | FunctionReference -> Call +!PARSE-TREE: | | | | | | ProcedureDesignator -> Name = 'mymax' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_out' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_in' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> Name = 'omp_in' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=0._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '0_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$omp declare reduction(min:tt:omp_out%r = min(omp_out%r, omp_in%r)) initializer(omp_priv%r = 1) -!CHECK-NEXT: !$OMP DECLARE REDUCTION(min:tt: omp_out%r = min(omp_out%r,omp_in%r)) INITIALIZER(omp_priv%r = 1_4) +!CHECK-NEXT: !$OMP DECLARE REDUCTION(min:tt: omp_out%r = min(omp_out%r, omp_in%r)) INITIALI& +!CHECK-NEXT: !$OMP&ZER(omp_priv%r = 1) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -70,8 +152,38 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'min' !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%r=min(omp_out%r,omp_in%r)' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=1._4' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out%r=min(omp_out%r,omp_in%r)' +!PARSE-TREE: | | | | Variable = 'omp_out%r' +!PARSE-TREE: | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Name = 'r' +!PARSE-TREE: | | | | Expr = 'min(omp_out%r,omp_in%r)' +!PARSE-TREE: | | | | | FunctionReference -> Call +!PARSE-TREE: | | | | | | ProcedureDesignator -> Name = 'min' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_out%r' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | Name = 'r' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_in%r' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | Name = 'r' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=1._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '1_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '1' +!PARSE-TREE: | Flags = None call random_number(values%r) diff --git a/flang/test/Parser/OpenMP/declare-reduction-operator.f90 b/flang/test/Parser/OpenMP/declare-reduction-operator.f90 index e4d07c8..0d337c1 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-operator.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-operator.f90 @@ -16,7 +16,8 @@ subroutine reduce_1 ( n, tts ) type(tt) :: tts(n) type(tt2) :: tts2(n) -!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out = tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)) INITIALIZER(omp_priv = tt(x=0_4,y=0_4)) +!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out = tt(omp_out%x - omp_in%x , omp_out%y - & +!CHECK: !$OMP&omp_in%y)) INITIALIZER(omp_priv = tt(0,0)) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -24,13 +25,60 @@ subroutine reduce_1 ( n, tts ) !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv=tt(x=0_4,y=0_4)' - +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | | StructureConstructor +!PARSE-TREE: | | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | | Name = 'tt' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%x-omp_in%x' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%y-omp_in%y' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv=tt(x=0_4,y=0_4)' +!PARSE-TREE: | | | Variable = 'omp_priv' +!PARSE-TREE: | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | Expr = 'tt(x=0_4,y=0_4)' +!PARSE-TREE: | | | | StructureConstructor +!PARSE-TREE: | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | Name = 'tt' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$omp declare reduction(+ : tt : omp_out = tt(omp_out%x - omp_in%x , omp_out%y - omp_in%y)) initializer(omp_priv = tt(0,0)) -!CHECK: !$OMP DECLARE REDUCTION(+:tt2: omp_out = tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)) INITIALIZER(omp_priv = tt2(x=0._8,y=0._8) +!CHECK: !$OMP DECLARE REDUCTION(+:tt2: omp_out = tt2(omp_out%x - omp_in%x , omp_out%y & +!CHECK: !$OMP&- omp_in%y)) INITIALIZER(omp_priv = tt2(0,0)) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -38,9 +86,55 @@ subroutine reduce_1 ( n, tts ) !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt2' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv=tt2(x=0._8,y=0._8)' - +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | | StructureConstructor +!PARSE-TREE: | | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | | Name = 'tt2' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%x-omp_in%x' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%y-omp_in%y' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv=tt2(x=0._8,y=0._8)' +!PARSE-TREE: | | | Variable = 'omp_priv' +!PARSE-TREE: | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | Expr = 'tt2(x=0._8,y=0._8)' +!PARSE-TREE: | | | | StructureConstructor +!PARSE-TREE: | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | Name = 'tt2' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$omp declare reduction(+ :tt2 : omp_out = tt2(omp_out%x - omp_in%x , omp_out%y - omp_in%y)) initializer(omp_priv = tt2(0,0)) type(tt) :: diffp = tt( 0, 0 ) diff --git a/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 b/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 index 455fc17..f026f15 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 @@ -8,6 +8,6 @@ end !CHECK: !DEF: /f00 (Subroutine) Subprogram !CHECK: subroutine f00 -!CHECK: !$omp declare reduction(fred:integer,real: omp_out = omp_in+omp_out) +!CHECK: !$omp declare reduction(fred:integer, real: omp_out = omp_in + omp_out) !CHECK: end subroutine diff --git a/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 b/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 index 73d7ccf..7897eb0 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 @@ -19,7 +19,8 @@ function func(x, n, init) end subroutine initme end interface !$omp declare reduction(red_add:integer(4):omp_out=omp_out+omp_in) initializer(initme(omp_priv,0)) -!CHECK: !$OMP DECLARE REDUCTION(red_add:INTEGER(KIND=4_4): omp_out = omp_out+omp_in) INITIALIZER(initme(omp_priv, 0_4)) +!CHECK: !$OMP DECLARE REDUCTION(red_add:INTEGER(KIND=4_4): omp_out=omp_out+omp_in) INITIA& +!CHECKL !$OMP&LIZER(initme(omp_priv,0)) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -27,9 +28,31 @@ function func(x, n, init) !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'red_add' !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec -> KindSelector -> Scalar -> Integer -> Constant -> Expr = '4_4' !PARSE-TREE: | | | LiteralConstant -> IntLiteralConstant = '4' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=omp_out+omp_in' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerProc -!PARSE-TREE: | | ProcedureDesignator -> Name = 'initme' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=omp_out+omp_in' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'omp_out+omp_in' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Expr = 'omp_in' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_in' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> CallStmt = 'CALL initme(omp_priv,0_4)' +!PARSE-TREE: | | | Call +!PARSE-TREE: | | | | ProcedureDesignator -> Name = 'initme' +!PARSE-TREE: | | | | ActualArgSpec +!PARSE-TREE: | | | | | ActualArg -> Expr = 'omp_priv' +!PARSE-TREE: | | | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | ActualArgSpec +!PARSE-TREE: | | | | | ActualArg -> Expr = '0_4' +!PARSE-TREE: | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None res=init !$omp simd reduction(red_add:res) @@ -59,7 +82,8 @@ end function func !CHECK-LABEL: program main program main integer :: my_var -!CHECK: !$OMP DECLARE REDUCTION(my_add_red:INTEGER: omp_out = omp_out+omp_in) INITIALIZER(omp_priv = 0_4) +!CHECK: !$OMP DECLARE REDUCTION(my_add_red:INTEGER: omp_out = omp_out + omp_in) INITIA& +!CHECK: !$OMP&LIZER(omp_priv=0) !$omp declare reduction (my_add_red : integer : omp_out = omp_out + omp_in) initializer (omp_priv=0) my_var = 0 @@ -74,5 +98,24 @@ end program main !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'my_add_red' !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec -> -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=omp_out+omp_in' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv=0_4' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=omp_out+omp_in' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'omp_out+omp_in' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Expr = 'omp_in' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_in' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv=0_4' +!PARSE-TREE: | | | Variable = 'omp_priv' +!PARSE-TREE: | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | Expr = '0_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None diff --git a/flang/test/Parser/OpenMP/metadirective-dirspec.f90 b/flang/test/Parser/OpenMP/metadirective-dirspec.f90 index c373001..b64ceb1 100644 --- a/flang/test/Parser/OpenMP/metadirective-dirspec.f90 +++ b/flang/test/Parser/OpenMP/metadirective-dirspec.f90 @@ -105,8 +105,8 @@ end !UNPARSE: TYPE :: tt2 !UNPARSE: REAL :: x !UNPARSE: END TYPE -!UNPARSE: !$OMP METADIRECTIVE WHEN(USER={CONDITION(.true._4)}: DECLARE REDUCTION(+:tt1,tt2: omp_out%x = omp_in%x+omp_out%x)& -!UNPARSE: !$OMP&) +!UNPARSE: !$OMP METADIRECTIVE WHEN(USER={CONDITION(.true._4)}: DECLARE REDUCTION(+:tt1, tt2: omp& +!UNPARSE: !$OMP&_out%x = omp_in%x + omp_out%x)) !UNPARSE: END SUBROUTINE !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OmpMetadirectiveDirective @@ -127,21 +127,44 @@ end !PARSE-TREE: | | | | | Name = 'tt1' !PARSE-TREE: | | | | OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | | | Name = 'tt2' -!PARSE-TREE: | | | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%x=omp_in%x+omp_out%x' -!PARSE-TREE: | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | DataRef -> Name = 'omp_out' -!PARSE-TREE: | | | | | | | Name = 'x' -!PARSE-TREE: | | | | | Expr = 'omp_in%x+omp_out%x' -!PARSE-TREE: | | | | | | Add -!PARSE-TREE: | | | | | | | Expr = 'omp_in%x' -!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_in' -!PARSE-TREE: | | | | | | | | | Name = 'x' -!PARSE-TREE: | | | | | | | Expr = 'omp_out%x' -!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_out' -!PARSE-TREE: | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | Instance -> AssignmentStmt = 'omp_out%x=omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | Variable = 'omp_out%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | Expr = 'omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | | Add +!PARSE-TREE: | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | OmpStylizedInstance +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | Instance -> AssignmentStmt = 'omp_out%x=omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | Variable = 'omp_out%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | Expr = 'omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | | Add +!PARSE-TREE: | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | Name = 'x' !PARSE-TREE: | | | OmpClauseList -> +!PARSE-TREE: | | | Flags = None subroutine f04 !$omp metadirective when(user={condition(.true.)}: & diff --git a/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 b/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 index 39e8f05..50a38c6 100644 --- a/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 +++ b/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 @@ -79,7 +79,7 @@ end !UNPARSE: TYPE :: t !UNPARSE: INTEGER :: x !UNPARSE: END TYPE -!UNPARSE: !$OMP DECLARE_REDUCTION(+:t: omp_out%x = omp_out%x+omp_in%x) +!UNPARSE: !$OMP DECLARE_REDUCTION(+:t: omp_out%x = omp_out%x + omp_in%x) !UNPARSE: END SUBROUTINE !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification @@ -88,21 +88,24 @@ end !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 't' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%x=omp_out%x+omp_in%x' -!PARSE-TREE: | | | Variable = 'omp_out%x' -!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | DataRef -> Name = 'omp_out' -!PARSE-TREE: | | | | | Name = 'x' -!PARSE-TREE: | | | Expr = 'omp_out%x+omp_in%x' -!PARSE-TREE: | | | | Add -!PARSE-TREE: | | | | | Expr = 'omp_out%x' -!PARSE-TREE: | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | DataRef -> Name = 'omp_out' -!PARSE-TREE: | | | | | | | Name = 'x' -!PARSE-TREE: | | | | | Expr = 'omp_in%x' -!PARSE-TREE: | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | DataRef -> Name = 'omp_in' -!PARSE-TREE: | | | | | | | Name = 'x' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out%x=omp_out%x+omp_in%x' +!PARSE-TREE: | | | | Variable = 'omp_out%x' +!PARSE-TREE: | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Name = 'x' +!PARSE-TREE: | | | | Expr = 'omp_out%x+omp_in%x' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | Name = 'x' !PARSE-TREE: | OmpClauseList -> !PARSE-TREE: | Flags = None diff --git a/flang/test/Semantics/OpenMP/declare-reduction-error.f90 b/flang/test/Semantics/OpenMP/declare-reduction-error.f90 deleted file mode 100644 index 21f5cc1..0000000 --- a/flang/test/Semantics/OpenMP/declare-reduction-error.f90 +++ /dev/null @@ -1,11 +0,0 @@ -! RUN: not %flang_fc1 -emit-obj -fopenmp -fopenmp-version=50 %s 2>&1 | FileCheck %s - -subroutine initme(x,n) - integer x,n - x=n -end subroutine initme - -subroutine subr - !$omp declare reduction(red_add:integer(4):omp_out=omp_out+omp_in) initializer(initme(omp_priv,0)) - !CHECK: error: Implicit subroutine declaration 'initme' in DECLARE REDUCTION -end subroutine subr diff --git a/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 b/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 index 000d323..89e0771 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 @@ -57,9 +57,10 @@ contains !CHECK: adder: UserReductionDetails TYPE(two) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !$omp simd reduction(adder:res) @@ -101,14 +102,16 @@ contains !CHECK: adder: UserReductionDetails TYPE(two) TYPE(three) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !CHECK OtherConstruct scope !CHECK: omp_in size=24 offset=0: ObjectEntity type: TYPE(three) -!CHECK: omp_orig size=24 offset=24: ObjectEntity type: TYPE(three) -!CHECK: omp_out size=24 offset=48: ObjectEntity type: TYPE(three) -!CHECK: omp_priv size=24 offset=72: ObjectEntity type: TYPE(three) +!CHECK: omp_out size=24 offset=24: ObjectEntity type: TYPE(three) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=24 offset=0: ObjectEntity type: TYPE(three) +!CHECK: omp_priv size=24 offset=24: ObjectEntity type: TYPE(three) !$omp simd reduction(adder:res3) do i=1,n @@ -135,9 +138,10 @@ contains !CHECK: op.+: UserReductionDetails TYPE(two) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !$omp simd reduction(+:res) @@ -163,14 +167,16 @@ contains !CHECK: op.+: UserReductionDetails TYPE(two) TYPE(three) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !CHECK: OtherConstruct scope !CHECK: omp_in size=24 offset=0: ObjectEntity type: TYPE(three) -!CHECK: omp_orig size=24 offset=24: ObjectEntity type: TYPE(three) -!CHECK: omp_out size=24 offset=48: ObjectEntity type: TYPE(three) -!CHECK: omp_priv size=24 offset=72: ObjectEntity type: TYPE(three) +!CHECK: omp_out size=24 offset=24: ObjectEntity type: TYPE(three) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=24 offset=0: ObjectEntity type: TYPE(three) +!CHECK: omp_priv size=24 offset=24: ObjectEntity type: TYPE(three) !$omp simd reduction(+:res3) do i=1,n @@ -183,6 +189,7 @@ contains enddo res%t2 = res2 res%t3 = res3 + funcBtwothree = res end function funcBtwothree !! This is checking a special case, where a reduction is declared inside a @@ -191,11 +198,12 @@ contains pure logical function reduction() !CHECK: reduction size=4 offset=0: ObjectEntity funcResult type: LOGICAL(4) !CHECK: rr: UserReductionDetails INTEGER(4) -!CHECK: OtherConstruct scope: size=16 alignment=4 sourceRange=0 bytes +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes !CHECK: omp_in size=4 offset=0: ObjectEntity type: INTEGER(4) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: INTEGER(4) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: INTEGER(4) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: INTEGER(4) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: INTEGER(4) +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: INTEGER(4) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: INTEGER(4) !$omp declare reduction (rr : integer : omp_out = omp_out + omp_in) initializer (omp_priv = 0) reduction = .false. end function reduction diff --git a/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 b/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 index 7ab7cad..87fcecd 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 @@ -18,9 +18,10 @@ contains !CHECK: op.AND: UserReductionDetails TYPE(logicalwrapper) !CHECK OtherConstruct scope !CHECK: omp_in size=4 offset=0: ObjectEntity type: TYPE(logicalwrapper) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: TYPE(logicalwrapper) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: TYPE(logicalwrapper) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: TYPE(logicalwrapper) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: TYPE(logicalwrapper) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: TYPE(logicalwrapper) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: TYPE(logicalwrapper) !$omp simd reduction(.AND.:res) do i=1,n diff --git a/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 b/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 index 0882de8..763179c 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 @@ -6,13 +6,13 @@ !type::t1 !integer(4)::val !endtype -!!$OMP DECLARE REDUCTION(*:t1:omp_out=omp_out*omp_in)INITIALIZER(omp_priv=& -!!$OMP&t1(1)) +!!$OMP DECLARE REDUCTION(*:t1: omp_out=omp_out*omp_in) INITIALIZER(omp_priv=t1(& +!!$OMP&1)) !!$OMP METADIRECTIVE OTHERWISE(DECLARE REDUCTION(+:INTEGER)) -!!$OMP DECLARE REDUCTION(.fluffy.:t1:omp_out=omp_out.fluffy.omp_in)INITIALI& -!!$OMP&ZER(omp_priv=t1(0)) -!!$OMP DECLARE REDUCTION(.mul.:t1:omp_out=omp_out.mul.omp_in)INITIALIZER(om& -!!$OMP&p_priv=t1(1)) +!!$OMP DECLARE REDUCTION(.fluffy.:t1: omp_out=omp_out.fluffy.omp_in) INITIALIZE& +!!$OMP&R(omp_priv=t1(0)) +!!$OMP DECLARE REDUCTION(.mul.:t1: omp_out=omp_out.mul.omp_in) INITIALIZER(omp_& +!!$OMP&priv=t1(1)) !interface operator(.mul.) !procedure::mul !end interface diff --git a/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 b/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 index dc12332..5fc4205 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 @@ -11,11 +11,9 @@ module m1 !$omp declare reduction(.fluffy.:t1:omp_out=omp_out.fluffy.omp_in) !CHECK: op.fluffy., PUBLIC: UserReductionDetails TYPE(t1) !CHECK: t1, PUBLIC: DerivedType components: val -!CHECK: OtherConstruct scope: size=16 alignment=4 sourceRange=0 bytes +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes !CHECK: omp_in size=4 offset=0: ObjectEntity type: TYPE(t1) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: TYPE(t1) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: TYPE(t1) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: TYPE(t1) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: TYPE(t1) contains function my_mul(x, y) type (t1), intent (in) :: x, y diff --git a/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 b/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 index 84dbe1a..e0006bf 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 @@ -64,9 +64,10 @@ program test_vector !CHECK: OtherConstruct scope: !CHECK: omp_in size=12 offset=0: ObjectEntity type: TYPE(vector) -!CHECK: omp_orig size=12 offset=12: ObjectEntity type: TYPE(vector) -!CHECK: omp_out size=12 offset=24: ObjectEntity type: TYPE(vector) -!CHECK: omp_priv size=12 offset=36: ObjectEntity type: TYPE(vector) +!CHECK: omp_out size=12 offset=12: ObjectEntity type: TYPE(vector) +!CHECK: OtherConstruct scope: +!CHECK: omp_orig size=12 offset=0: ObjectEntity type: TYPE(vector) +!CHECK: omp_priv size=12 offset=12: ObjectEntity type: TYPE(vector) v2 = Vector(0.0, 0.0, 0.0) v1 = Vector(1.0, 2.0, 3.0) diff --git a/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 b/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 index 9cd638d..115fe51 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 @@ -33,11 +33,12 @@ program test_omp_reduction !$omp declare reduction (.modmul. : t1 : omp_out = omp_out .modmul. omp_in) initializer(omp_priv = t1(1.0)) !CHECK: op.modmul.: UserReductionDetails TYPE(t1) !CHECK: t1: Use from t1 in module1 -!CHECK: OtherConstruct scope: size=16 alignment=4 sourceRange=0 bytes +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes !CHECK: omp_in size=4 offset=0: ObjectEntity type: TYPE(t1) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: TYPE(t1) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: TYPE(t1) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: TYPE(t1) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: TYPE(t1) +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: TYPE(t1) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: TYPE(t1) result = t1(1.0) !$omp parallel do reduction(.modmul.:result) do i = 1, 10 diff --git a/flang/test/Semantics/OpenMP/declare-reduction.f90 b/flang/test/Semantics/OpenMP/declare-reduction.f90 index 1f39c57..c8dee5e 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction.f90 @@ -19,10 +19,12 @@ function func(x, n, init) !$omp declare reduction(red_add:integer(4):omp_out=omp_out+omp_in) initializer(initme(omp_priv,0)) !CHECK: red_add: UserReductionDetails !CHECK: Subprogram scope: initme +!CHECK: OtherConstruct scope: !CHECK: omp_in size=4 offset=0: ObjectEntity type: INTEGER(4) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: INTEGER(4) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: INTEGER(4) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: INTEGER(4) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: INTEGER(4) +!CHECK: OtherConstruct scope: +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: INTEGER(4) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: INTEGER(4) !$omp simd reduction(red_add:res) do i=1,n res=res+x(i) @@ -36,9 +38,11 @@ program main !$omp declare reduction (my_add_red : integer : omp_out = omp_out + omp_in) initializer (omp_priv=0) !CHECK: my_add_red: UserReductionDetails +!CHECK: OtherConstruct scope: !CHECK: omp_in size=4 offset=0: ObjectEntity type: INTEGER(4) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: INTEGER(4) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: INTEGER(4) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: INTEGER(4) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: INTEGER(4) +!CHECK: OtherConstruct scope: +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: INTEGER(4) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: INTEGER(4) end program main diff --git a/flang/test/Semantics/ignore_tkr04.f90 b/flang/test/Semantics/ignore_tkr04.f90 new file mode 100644 index 0000000..8becc85 --- /dev/null +++ b/flang/test/Semantics/ignore_tkr04.f90 @@ -0,0 +1,26 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 +! Tests for ignore_tkr(p) +module ignore_tkr_4_m +interface + subroutine s(a) + real, pointer :: a(:) +!dir$ ignore_tkr(p) a + end subroutine + subroutine s1(a) + real, allocatable :: a(:) +!dir$ ignore_tkr(p) a + end subroutine +end interface +end module +program t + use ignore_tkr_4_m + real, allocatable :: x(:) + real, pointer :: x1(:) + call s(x) +!CHECK-NOT: error +!CHECK-NOT: warning + call s1(x1) +!CHECK-NOT: error +!CHECK-NOT: warning +end + diff --git a/flang/unittests/CMakeLists.txt b/flang/unittests/CMakeLists.txt index db04923e..2d612e58 100644 --- a/flang/unittests/CMakeLists.txt +++ b/flang/unittests/CMakeLists.txt @@ -48,6 +48,7 @@ function(add_flang_nongtest_unittest test_name) llvm_map_components_to_libnames(llvm_libs Support) endif() target_link_libraries(${test_name}${suffix} ${llvm_libs} ${ARG_UNPARSED_ARGUMENTS}) + set_unittest_link_flags(${test_name}${suffix}) if(NOT ARG_SLOW_TEST) add_dependencies(FlangUnitTests ${test_name}${suffix}) |