aboutsummaryrefslogtreecommitdiff
path: root/flang/lib
diff options
context:
space:
mode:
Diffstat (limited to 'flang/lib')
-rw-r--r--flang/lib/Optimizer/Builder/IntrinsicCall.cpp315
-rw-r--r--flang/lib/Optimizer/CodeGen/TargetRewrite.cpp9
-rw-r--r--flang/lib/Optimizer/Dialect/FIROps.cpp7
-rw-r--r--flang/lib/Parser/openmp-parsers.cpp244
-rw-r--r--flang/lib/Parser/openmp-utils.cpp12
-rw-r--r--flang/lib/Parser/parse-tree.cpp27
-rw-r--r--flang/lib/Parser/unparse.cpp37
-rw-r--r--flang/lib/Semantics/resolve-directives.cpp17
-rw-r--r--flang/lib/Semantics/resolve-names.cpp99
9 files changed, 659 insertions, 108 deletions
diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index 39bac81..53fe9c0 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,93 @@ 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);
+ 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 +9451,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> &regions) {
// 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/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..0e6d4c7 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));
}
generated by cgit v1.1 at 2025-10-29 14:51:40 +0000