[OpenACC][CIR] Implement 'device_type' clause lowering for 'init'/'sh… #135102
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…utdown' 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.
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) Changes…utdown' 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 Full diff: https://github.com/llvm/llvm-project/pull/135102.diff 4 Files Affected:
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>]}
}
|
|
@llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) Changes…utdown' 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 Full diff: https://github.com/llvm/llvm-project/pull/135102.diff 4 Files Affected:
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>]}
}
|
clementval
approved these changes
Apr 10, 2025
| @@ -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(*) | |||
There was a problem hiding this comment.
What happens if you have this?
#pragma acc init device_type(*) device_type(nvidia)
There was a problem hiding this comment.
This will faithfully represent them, like this:
// CHECK-NEXT: acc.init attributes {device_types = [#acc.device_type<star>, #acc.device_type<nvidia>]}
There isn't really anything that prohibits it by standard, so it seems reasonable to do.
andykaylor
reviewed
Apr 10, 2025
| @@ -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, | |||
There was a problem hiding this comment.
I don't like how inconsistent we've been about the placement of the location parameter. I thought there was a tendency in MLIR for it to be the first argument (though that felt odd to me). I just looked at the emit functions in the incubator and it's all over the place.
Do you think there's value in standardizing on it being the first argument?
| // '*' case leaves no identifier-info, just a nullptr. | ||
| if (!II) | ||
| return mlir::acc::DeviceType::Star; | ||
| return llvm::StringSwitch<mlir::acc::DeviceType>(II->getName()) |
There was a problem hiding this comment.
What does this return if none of the strings is matched?
There was a problem hiding this comment.
It doesn't. It asserts, which is the behavior we want.
| @@ -57,31 +75,90 @@ class OpenACCClauseCIREmitter final | |||
| } | |||
| } | |||
|
|
|||
| mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *II) { | |||
|
|
|||
There was a problem hiding this comment.
I've been telling people not to put a blank line here because it seemed inconsistent with the prevailing style in clang. Am I wrong about that?
There was a problem hiding this comment.
I'm not sure there is a style guidance here. But I'll remove it as it is probably wrong.
| @@ -57,31 +75,90 @@ class OpenACCClauseCIREmitter final | |||
| } | |||
| } | |||
|
|
|||
| mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *II) { | |||
There was a problem hiding this comment.
Suggested change
| mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *II) { | |
| mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) { |
I don't think we make exceptions for acronyms.
| // 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()) |
There was a problem hiding this comment.
Suggested change
| for (const DeviceTypeArgument &Arg : clause.getArchitectures()) | |
| for (const DeviceTypeArgument &arg : clause.getArchitectures()) |
| 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>) |
There was a problem hiding this comment.
Am I right that ParallelOp is acc::ParallelOp?
I was confused about what was happening here because I didn't read the comment where isOneOfType was defined. A brief comment here explaining what you're looking for ('Is Op one of the expected types?") would be helpful.
There was a problem hiding this comment.
yep, ParallelOp is mlir::acc::ParallelOp. We have a using namespace mlir::acc above. I'll add the comment.
| op.setDeviceTypesAttr( | ||
| mlir::ArrayAttr::get(builder.getContext(), deviceTypes)); | ||
| } else { | ||
| cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ", |
There was a problem hiding this comment.
Is this really "NYI"? It feels more like we shouldn't get here. In either case, since the condition is static, can this be a static compile error?
There was a problem hiding this comment.
It is, there are entries for Device_type that aren't implemented yet, the list on 126 is incomplete. Eventually, this will become a no-op/unreachable.
It can't be a static_assert or anything because the condition on line 122 is not constexpr. So we would still have to evaluate the static_assert in every other case.
| @@ -96,18 +173,21 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt( | |||
|
|
|||
| template <typename Op> | |||
| mlir::LogicalResult | |||
| CIRGenFunction::emitOpenACCOp(mlir::Location start, | |||
| CIRGenFunction::emitOpenACCOp(OpenACCDirectiveKind dirKind, | |||
| SourceLocation dirLoc, mlir::Location start, | |||
There was a problem hiding this comment.
Why not use mlir::Location for dirLoc?
There was a problem hiding this comment.
Because that location is being used for NYI diagnostics, which take the Clang Source Location, not the mlir Location.
There was a problem hiding this comment.
I see. It's unfortunate to have a parameter that's just there for NYI diagnostics. I assume you plan to remove it when everything is implemented?
There was a problem hiding this comment.
Yep, exactly my intent. I struggled a bit deciding how I could do the NYI without it, but figure I'll remove it once everythign is implemented.
var-const
pushed a commit
to ldionne/llvm-project
that referenced
this pull request
Apr 17, 2025
llvm#135102) …utdown' 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.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
…utdown'
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
applyAttributesfunction 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 someif constexprand a small type-trait to help legalize these.