https://github.com/erichkeane updated
https://github.com/llvm/llvm-project/pull/135102
>From e7ee7372e9e5df14b6aa893005e842eaa2ae49ff Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Wed, 9 Apr 2025 16:35:06 -0700
Subject: [PATCH 1/2] [OpenACC][CIR] Implement 'device_type' clause lowering
for 'init'/'shutdown'
This patch emits the lowering for 'device_type' on an 'init' or
'shutdown'. This one is fairly unique, as these directives have it as an
attribute, rather than as a component of the individual operands, like
the rest of the constructs.
So this patch implements the lowering as an attribute.
In order to do tis, a few refactorings had to happen:
First, the 'emitOpenACCOp' functions needed to pick up th edirective
kind/location so that the NYI diagnostic could be reasonable.
Second, and most impactful, the `applyAttributes` function ends up
needing to encode some of the appertainment rules, thanks to the way the
OpenACC-MLIR operands get their attributes attached. Since they each
use a special function (rather than something that can be legalized at
runtime), the forms of 'setDefaultAttr' is only valid for some ops. SO
this patch uses some `if constexpr` and a small type-trait to help
legalize these.
---
clang/lib/CIR/CodeGen/CIRGenFunction.h | 11 +-
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 120 +++++++++++++++++---
clang/test/CIR/CodeGenOpenACC/init.c | 13 +++
clang/test/CIR/CodeGenOpenACC/shutdown.c | 13 +++
4 files changed, 135 insertions(+), 22 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h
b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 6ffa106f2a383..53b072fbba00f 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -585,15 +585,16 @@ class CIRGenFunction : public CIRGenTypeCache {
private:
template <typename Op>
mlir::LogicalResult
- emitOpenACCOp(mlir::Location start,
+ emitOpenACCOp(OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
+ mlir::Location start,
llvm::ArrayRef<const OpenACCClause *> clauses);
// Function to do the basic implementation of an operation with an Associated
// Statement. Models AssociatedStmtConstruct.
template <typename Op, typename TermOp>
- mlir::LogicalResult
- emitOpenACCOpAssociatedStmt(mlir::Location start, mlir::Location end,
- llvm::ArrayRef<const OpenACCClause *> clauses,
- const Stmt *associatedStmt);
+ mlir::LogicalResult emitOpenACCOpAssociatedStmt(
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc, mlir::Location
start,
+ mlir::Location end, llvm::ArrayRef<const OpenACCClause *> clauses,
+ const Stmt *associatedStmt);
public:
mlir::LogicalResult
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 8c3c87a58c269..b4c887945461b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -9,6 +9,7 @@
// Emit OpenACC Stmt nodes as CIR code.
//
//===----------------------------------------------------------------------===//
+#include <type_traits>
#include "CIRGenBuilder.h"
#include "CIRGenFunction.h"
@@ -23,14 +24,29 @@ using namespace cir;
using namespace mlir::acc;
namespace {
+// Simple type-trait to see if the first template arg is one of the list, so we
+// can tell whether to `if-constexpr` a bunch of stuff.
+template <typename ToTest, typename T, typename... Tys>
+constexpr bool isOneOfTypes =
+ std::is_same_v<ToTest, T> || isOneOfTypes<ToTest, Tys...>;
+template <typename ToTest, typename T>
+constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
+
class OpenACCClauseCIREmitter final
: public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
CIRGenModule &cgm;
+ // This is necessary since a few of the clauses emit differently based on the
+ // directive kind they are attached to.
+ OpenACCDirectiveKind dirKind;
+ SourceLocation dirLoc;
struct AttributeData {
// Value of the 'default' attribute, added on 'data' and 'compute'/etc
// constructs as a 'default-attr'.
std::optional<ClauseDefaultValue> defaultVal = std::nullopt;
+ // For directives that have their device type architectures listed in
+ // attributes (init/shutdown/etc), the list of architectures to be emitted.
+ llvm::SmallVector<mlir::acc::DeviceType> deviceTypeArchs{};
} attrData;
void clauseNotImplemented(const OpenACCClause &c) {
@@ -38,7 +54,9 @@ class OpenACCClauseCIREmitter final
}
public:
- OpenACCClauseCIREmitter(CIRGenModule &cgm) : cgm(cgm) {}
+ OpenACCClauseCIREmitter(CIRGenModule &cgm, OpenACCDirectiveKind dirKind,
+ SourceLocation dirLoc)
+ : cgm(cgm), dirKind(dirKind), dirLoc(dirLoc) {}
void VisitClause(const OpenACCClause &clause) {
clauseNotImplemented(clause);
@@ -57,31 +75,90 @@ class OpenACCClauseCIREmitter final
}
}
+ mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *II) {
+
+ // '*' case leaves no identifier-info, just a nullptr.
+ if (!II)
+ return mlir::acc::DeviceType::Star;
+ return llvm::StringSwitch<mlir::acc::DeviceType>(II->getName())
+ .CaseLower("default", mlir::acc::DeviceType::Default)
+ .CaseLower("host", mlir::acc::DeviceType::Host)
+ .CaseLower("multicore", mlir::acc::DeviceType::Multicore)
+ .CasesLower("nvidia", "acc_device_nvidia",
+ mlir::acc::DeviceType::Nvidia)
+ .CaseLower("radeon", mlir::acc::DeviceType::Radeon);
+ }
+
+ void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
+
+ switch (dirKind) {
+ case OpenACCDirectiveKind::Init:
+ case OpenACCDirectiveKind::Shutdown: {
+ // Device type has a list that is either a 'star' (emitted as 'star'),
+ // or an identifer list, all of which get added for attributes.
+
+ for (const DeviceTypeArgument &Arg : clause.getArchitectures())
+ attrData.deviceTypeArchs.push_back(decodeDeviceType(Arg.first));
+ break;
+ }
+ default:
+ return clauseNotImplemented(clause);
+ }
+ }
+
// Apply any of the clauses that resulted in an 'attribute'.
- template <typename Op> void applyAttributes(Op &op) {
- if (attrData.defaultVal.has_value())
- op.setDefaultAttr(*attrData.defaultVal);
+ template <typename Op>
+ void applyAttributes(CIRGenBuilderTy &builder, Op &op) {
+
+ if (attrData.defaultVal.has_value()) {
+ // FIXME: OpenACC: as we implement this for other directive kinds, we
have
+ // to expand this list.
+ if constexpr (isOneOfTypes<Op, ParallelOp, SerialOp, KernelsOp, DataOp>)
+ op.setDefaultAttr(*attrData.defaultVal);
+ else
+ cgm.errorNYI(dirLoc, "OpenACC 'default' clause lowering for ",
dirKind);
+ }
+
+ if (!attrData.deviceTypeArchs.empty()) {
+ // FIXME: OpenACC: as we implement this for other directive kinds, we
have
+ // to expand this list, or more likely, have a 'noop' branch as most
other
+ // uses of this apply to the operands instead.
+ if constexpr (isOneOfTypes<Op, InitOp, ShutdownOp>) {
+ llvm::SmallVector<mlir::Attribute> deviceTypes;
+ for (mlir::acc::DeviceType DT : attrData.deviceTypeArchs)
+ deviceTypes.push_back(
+ mlir::acc::DeviceTypeAttr::get(builder.getContext(), DT));
+
+ op.setDeviceTypesAttr(
+ mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
+ } else {
+ cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ",
+ dirKind);
+ }
+ }
}
};
+
} // namespace
template <typename Op, typename TermOp>
mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
- mlir::Location start, mlir::Location end,
- llvm::ArrayRef<const OpenACCClause *> clauses, const Stmt *associatedStmt)
{
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc, mlir::Location start,
+ mlir::Location end, llvm::ArrayRef<const OpenACCClause *> clauses,
+ const Stmt *associatedStmt) {
mlir::LogicalResult res = mlir::success();
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
// Clause-emitter must be here because it might modify operands.
- OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
+ OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc);
clauseEmitter.VisitClauseList(clauses);
auto op = builder.create<Op>(start, retTy, operands);
// Apply the attributes derived from the clauses.
- clauseEmitter.applyAttributes(op);
+ clauseEmitter.applyAttributes(builder, op);
mlir::Block &block = op.getRegion().emplaceBlock();
mlir::OpBuilder::InsertionGuard guardCase(builder);
@@ -96,7 +173,8 @@ mlir::LogicalResult
CIRGenFunction::emitOpenACCOpAssociatedStmt(
template <typename Op>
mlir::LogicalResult
-CIRGenFunction::emitOpenACCOp(mlir::Location start,
+CIRGenFunction::emitOpenACCOp(OpenACCDirectiveKind dirKind,
+ SourceLocation dirLoc, mlir::Location start,
llvm::ArrayRef<const OpenACCClause *> clauses) {
mlir::LogicalResult res = mlir::success();
@@ -104,10 +182,12 @@ CIRGenFunction::emitOpenACCOp(mlir::Location start,
llvm::SmallVector<mlir::Value> operands;
// Clause-emitter must be here because it might modify operands.
- OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
+ OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc);
clauseEmitter.VisitClauseList(clauses);
- builder.create<Op>(start, retTy, operands);
+ auto op = builder.create<Op>(start, retTy, operands);
+ // Apply the attributes derived from the clauses.
+ clauseEmitter.applyAttributes(builder, op);
return res;
}
@@ -119,13 +199,16 @@ CIRGenFunction::emitOpenACCComputeConstruct(const
OpenACCComputeConstruct &s) {
switch (s.getDirectiveKind()) {
case OpenACCDirectiveKind::Parallel:
return emitOpenACCOpAssociatedStmt<ParallelOp, mlir::acc::YieldOp>(
- start, end, s.clauses(), s.getStructuredBlock());
+ s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ s.getStructuredBlock());
case OpenACCDirectiveKind::Serial:
return emitOpenACCOpAssociatedStmt<SerialOp, mlir::acc::YieldOp>(
- start, end, s.clauses(), s.getStructuredBlock());
+ s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ s.getStructuredBlock());
case OpenACCDirectiveKind::Kernels:
return emitOpenACCOpAssociatedStmt<KernelsOp, mlir::acc::TerminatorOp>(
- start, end, s.clauses(), s.getStructuredBlock());
+ s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ s.getStructuredBlock());
default:
llvm_unreachable("invalid compute construct kind");
}
@@ -137,18 +220,21 @@ CIRGenFunction::emitOpenACCDataConstruct(const
OpenACCDataConstruct &s) {
mlir::Location end = getLoc(s.getSourceRange().getEnd());
return emitOpenACCOpAssociatedStmt<DataOp, mlir::acc::TerminatorOp>(
- start, end, s.clauses(), s.getStructuredBlock());
+ s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ s.getStructuredBlock());
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCInitConstruct(const OpenACCInitConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getEnd());
- return emitOpenACCOp<InitOp>(start, s.clauses());
+ return emitOpenACCOp<InitOp>(s.getDirectiveKind(), s.getDirectiveLoc(),
start,
+ s.clauses());
}
mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct(
const OpenACCShutdownConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getEnd());
- return emitOpenACCOp<ShutdownOp>(start, s.clauses());
+ return emitOpenACCOp<ShutdownOp>(s.getDirectiveKind(), s.getDirectiveLoc(),
+ start, s.clauses());
}
mlir::LogicalResult
diff --git a/clang/test/CIR/CodeGenOpenACC/init.c
b/clang/test/CIR/CodeGenOpenACC/init.c
index e81e211b2608f..38957ad7dce75 100644
--- a/clang/test/CIR/CodeGenOpenACC/init.c
+++ b/clang/test/CIR/CodeGenOpenACC/init.c
@@ -4,4 +4,17 @@ void acc_init(void) {
// CHECK: cir.func @acc_init() {
#pragma acc init
// CHECK-NEXT: acc.init loc(#{{[a-zA-Z0-9]+}}){{$}}
+
+#pragma acc init device_type(*)
+ // CHECK-NEXT: acc.init attributes {device_types = [#acc.device_type<star>]}
+#pragma acc init device_type(nvidia)
+ // CHECK-NEXT: acc.init attributes {device_types =
[#acc.device_type<nvidia>]}
+#pragma acc init device_type(host, multicore)
+ // CHECK-NEXT: acc.init attributes {device_types = [#acc.device_type<host>,
#acc.device_type<multicore>]}
+#pragma acc init device_type(NVIDIA)
+ // CHECK-NEXT: acc.init attributes {device_types =
[#acc.device_type<nvidia>]}
+#pragma acc init device_type(HoSt, MuLtIcORe)
+ // CHECK-NEXT: acc.init attributes {device_types = [#acc.device_type<host>,
#acc.device_type<multicore>]}
+#pragma acc init device_type(HoSt) device_type(MuLtIcORe)
+ // CHECK-NEXT: acc.init attributes {device_types = [#acc.device_type<host>,
#acc.device_type<multicore>]}
}
diff --git a/clang/test/CIR/CodeGenOpenACC/shutdown.c
b/clang/test/CIR/CodeGenOpenACC/shutdown.c
index f971807529ecd..c14e090b7edb7 100644
--- a/clang/test/CIR/CodeGenOpenACC/shutdown.c
+++ b/clang/test/CIR/CodeGenOpenACC/shutdown.c
@@ -4,4 +4,17 @@ void acc_shutdown(void) {
// CHECK: cir.func @acc_shutdown() {
#pragma acc shutdown
// CHECK-NEXT: acc.shutdown loc(#{{[a-zA-Z0-9]+}}){{$}}
+
+#pragma acc shutdown device_type(*)
+ // CHECK-NEXT: acc.shutdown attributes {device_types =
[#acc.device_type<star>]}
+#pragma acc shutdown device_type(nvidia)
+ // CHECK-NEXT: acc.shutdown attributes {device_types =
[#acc.device_type<nvidia>]}
+#pragma acc shutdown device_type(host, multicore)
+ // CHECK-NEXT: acc.shutdown attributes {device_types =
[#acc.device_type<host>, #acc.device_type<multicore>]}
+#pragma acc shutdown device_type(NVIDIA)
+ // CHECK-NEXT: acc.shutdown attributes {device_types =
[#acc.device_type<nvidia>]}
+#pragma acc shutdown device_type(HoSt, MuLtIcORe)
+ // CHECK-NEXT: acc.shutdown attributes {device_types =
[#acc.device_type<host>, #acc.device_type<multicore>]}
+#pragma acc shutdown device_type(HoSt) device_type(MuLtIcORe)
+ // CHECK-NEXT: acc.shutdown attributes {device_types =
[#acc.device_type<host>, #acc.device_type<multicore>]}
}
>From 7a1a753bddfa8c7f85f2f45bd2e3dbfd024597ad Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Thu, 10 Apr 2025 10:21:42 -0700
Subject: [PATCH 2/2] Fix things from andy's review
---
clang/lib/CIR/CodeGen/CIRGenFunction.h | 8 ++---
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 40 +++++++++++----------
2 files changed, 25 insertions(+), 23 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h
b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 53b072fbba00f..be202dbe76c63 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -585,15 +585,15 @@ class CIRGenFunction : public CIRGenTypeCache {
private:
template <typename Op>
mlir::LogicalResult
- emitOpenACCOp(OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
- mlir::Location start,
+ emitOpenACCOp(mlir::Location start, OpenACCDirectiveKind dirKind,
+ SourceLocation dirLoc,
llvm::ArrayRef<const OpenACCClause *> clauses);
// Function to do the basic implementation of an operation with an Associated
// Statement. Models AssociatedStmtConstruct.
template <typename Op, typename TermOp>
mlir::LogicalResult emitOpenACCOpAssociatedStmt(
- OpenACCDirectiveKind dirKind, SourceLocation dirLoc, mlir::Location
start,
- mlir::Location end, llvm::ArrayRef<const OpenACCClause *> clauses,
+ mlir::Location start, mlir::Location end, OpenACCDirectiveKind dirKind,
+ SourceLocation dirLoc, llvm::ArrayRef<const OpenACCClause *> clauses,
const Stmt *associatedStmt);
public:
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index b4c887945461b..6fc5a7e0dbd37 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -75,12 +75,11 @@ class OpenACCClauseCIREmitter final
}
}
- mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *II) {
-
+ mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) {
// '*' case leaves no identifier-info, just a nullptr.
- if (!II)
+ if (!ii)
return mlir::acc::DeviceType::Star;
- return llvm::StringSwitch<mlir::acc::DeviceType>(II->getName())
+ return llvm::StringSwitch<mlir::acc::DeviceType>(ii->getName())
.CaseLower("default", mlir::acc::DeviceType::Default)
.CaseLower("host", mlir::acc::DeviceType::Host)
.CaseLower("multicore", mlir::acc::DeviceType::Multicore)
@@ -97,8 +96,8 @@ class OpenACCClauseCIREmitter final
// Device type has a list that is either a 'star' (emitted as 'star'),
// or an identifer list, all of which get added for attributes.
- for (const DeviceTypeArgument &Arg : clause.getArchitectures())
- attrData.deviceTypeArchs.push_back(decodeDeviceType(Arg.first));
+ for (const DeviceTypeArgument &arg : clause.getArchitectures())
+ attrData.deviceTypeArchs.push_back(decodeDeviceType(arg.first));
break;
}
default:
@@ -113,6 +112,8 @@ class OpenACCClauseCIREmitter final
if (attrData.defaultVal.has_value()) {
// FIXME: OpenACC: as we implement this for other directive kinds, we
have
// to expand this list.
+ // This type-trait checks if 'op'(the first arg) is one of the mlir::acc
+ // operations listed in the rest of the arguments.
if constexpr (isOneOfTypes<Op, ParallelOp, SerialOp, KernelsOp, DataOp>)
op.setDefaultAttr(*attrData.defaultVal);
else
@@ -123,6 +124,7 @@ class OpenACCClauseCIREmitter final
// FIXME: OpenACC: as we implement this for other directive kinds, we
have
// to expand this list, or more likely, have a 'noop' branch as most
other
// uses of this apply to the operands instead.
+ // This type-trait checks if 'op'(the first arg) is one of the mlir::acc
if constexpr (isOneOfTypes<Op, InitOp, ShutdownOp>) {
llvm::SmallVector<mlir::Attribute> deviceTypes;
for (mlir::acc::DeviceType DT : attrData.deviceTypeArchs)
@@ -143,8 +145,8 @@ class OpenACCClauseCIREmitter final
template <typename Op, typename TermOp>
mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
- OpenACCDirectiveKind dirKind, SourceLocation dirLoc, mlir::Location start,
- mlir::Location end, llvm::ArrayRef<const OpenACCClause *> clauses,
+ mlir::Location start, mlir::Location end, OpenACCDirectiveKind dirKind,
+ SourceLocation dirLoc, llvm::ArrayRef<const OpenACCClause *> clauses,
const Stmt *associatedStmt) {
mlir::LogicalResult res = mlir::success();
@@ -172,10 +174,9 @@ mlir::LogicalResult
CIRGenFunction::emitOpenACCOpAssociatedStmt(
}
template <typename Op>
-mlir::LogicalResult
-CIRGenFunction::emitOpenACCOp(OpenACCDirectiveKind dirKind,
- SourceLocation dirLoc, mlir::Location start,
- llvm::ArrayRef<const OpenACCClause *> clauses) {
+mlir::LogicalResult CIRGenFunction::emitOpenACCOp(
+ mlir::Location start, OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
+ llvm::ArrayRef<const OpenACCClause *> clauses) {
mlir::LogicalResult res = mlir::success();
llvm::SmallVector<mlir::Type> retTy;
@@ -199,15 +200,15 @@ CIRGenFunction::emitOpenACCComputeConstruct(const
OpenACCComputeConstruct &s) {
switch (s.getDirectiveKind()) {
case OpenACCDirectiveKind::Parallel:
return emitOpenACCOpAssociatedStmt<ParallelOp, mlir::acc::YieldOp>(
- s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
s.getStructuredBlock());
case OpenACCDirectiveKind::Serial:
return emitOpenACCOpAssociatedStmt<SerialOp, mlir::acc::YieldOp>(
- s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
s.getStructuredBlock());
case OpenACCDirectiveKind::Kernels:
return emitOpenACCOpAssociatedStmt<KernelsOp, mlir::acc::TerminatorOp>(
- s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
s.getStructuredBlock());
default:
llvm_unreachable("invalid compute construct kind");
@@ -220,21 +221,22 @@ CIRGenFunction::emitOpenACCDataConstruct(const
OpenACCDataConstruct &s) {
mlir::Location end = getLoc(s.getSourceRange().getEnd());
return emitOpenACCOpAssociatedStmt<DataOp, mlir::acc::TerminatorOp>(
- s.getDirectiveKind(), s.getDirectiveLoc(), start, end, s.clauses(),
+ start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
s.getStructuredBlock());
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCInitConstruct(const OpenACCInitConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getEnd());
- return emitOpenACCOp<InitOp>(s.getDirectiveKind(), s.getDirectiveLoc(),
start,
+ return emitOpenACCOp<InitOp>(start, s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses());
}
+
mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct(
const OpenACCShutdownConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getEnd());
- return emitOpenACCOp<ShutdownOp>(s.getDirectiveKind(), s.getDirectiveLoc(),
- start, s.clauses());
+ return emitOpenACCOp<ShutdownOp>(start, s.getDirectiveKind(),
+ s.getDirectiveLoc(), s.clauses());
}
mlir::LogicalResult
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
