[OpenACC][CIR] Implement 'self' lowering on compute constructs #135851
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
This is our first attempt at lowering a clause that is an 'operand' in the OpenACC operand, so it does quite a bit of refactoring. My previous plans on how to emit the clauses was not viable, so we instead do 'create the op, then use the visitor to fill in the operands'. This resulted in the 'applyAttributes' function getting removed and a few other functions simplified. Additionally, it requires setting the insertion point a little to make sure we're inserting 'around' the operation correctly. Finally, since the OpenACC dialect only understands the MLIR types, we had to introduce a use of the unrealized-conversion-cast, which we'll probably getting good use out of in the future.
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThis is our first attempt at lowering a clause that is an 'operand' in the OpenACC operand, so it does quite a bit of refactoring. My previous plans on how to emit the clauses was not viable, so we instead do 'create the op, then use the visitor to fill in the operands'. This resulted in the 'applyAttributes' function getting removed and a few other functions simplified. Additionally, it requires setting the insertion point a little to make sure we're inserting 'around' the operation correctly. Finally, since the OpenACC dialect only understands the MLIR types, we had to introduce a use of the unrealized-conversion-cast, which we'll probably getting good use out of in the future. Full diff: https://github.com/llvm/llvm-project/pull/135851.diff 5 Files Affected:
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 3687af76a559f..681567228cbb0 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -430,6 +430,11 @@ class OpenACCSelfClause final
}
bool isConditionExprClause() const { return HasConditionExpr.has_value(); }
+ bool isVarListClause() const { return !isConditionExprClause(); }
+ bool isEmptySelfClause() const {
+ return (isConditionExprClause() && !hasConditionExpr()) ||
+ (!isConditionExprClause() && getVarList().empty());
+ }
bool hasConditionExpr() const {
assert(HasConditionExpr.has_value() &&
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 152f996ed0fed..92f3ad2a68eb8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -32,46 +32,52 @@ constexpr bool isOneOfTypes =
template <typename ToTest, typename T>
constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
+template <typename OpTy>
class OpenACCClauseCIREmitter final
- : public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
- CIRGenModule &cgm;
+ : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> {
+ OpTy &operation;
+ CIRGenFunction &cgf;
+ CIRGenBuilderTy &builder;
+
// This is necessary since a few of the clauses emit differently based on the
// directive kind they are attached to.
OpenACCDirectiveKind dirKind;
+ // This source location should be able to go away once the NYI diagnostics are
+ // gone.
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) {
- cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
+ cgf.getCIRGenModule().errorNYI(c.getSourceRange(), "OpenACC Clause",
+ c.getClauseKind());
}
public:
- OpenACCClauseCIREmitter(CIRGenModule &cgm, OpenACCDirectiveKind dirKind,
- SourceLocation dirLoc)
- : cgm(cgm), dirKind(dirKind), dirLoc(dirLoc) {}
+ OpenACCClauseCIREmitter(OpTy &operation, CIRGenFunction &cgf,
+ CIRGenBuilderTy &builder,
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc)
+ : operation(operation), cgf(cgf), builder(builder), dirKind(dirKind),
+ dirLoc(dirLoc) {}
void VisitClause(const OpenACCClause &clause) {
clauseNotImplemented(clause);
}
void VisitDefaultClause(const OpenACCDefaultClause &clause) {
- switch (clause.getDefaultClauseKind()) {
- case OpenACCDefaultClauseKind::None:
- attrData.defaultVal = ClauseDefaultValue::None;
- break;
- case OpenACCDefaultClauseKind::Present:
- attrData.defaultVal = ClauseDefaultValue::Present;
- break;
- case OpenACCDefaultClauseKind::Invalid:
- break;
+ // 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<OpTy, ParallelOp, SerialOp, KernelsOp, DataOp>) {
+ switch (clause.getDefaultClauseKind()) {
+ case OpenACCDefaultClauseKind::None:
+ operation.setDefaultAttr(ClauseDefaultValue::None);
+ break;
+ case OpenACCDefaultClauseKind::Present:
+ operation.setDefaultAttr(ClauseDefaultValue::Present);
+ break;
+ case OpenACCDefaultClauseKind::Invalid:
+ break;
+ }
+ } else {
+ return clauseNotImplemented(clause);
}
}
@@ -89,64 +95,70 @@ class OpenACCClauseCIREmitter final
}
void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) {
+ llvm::SmallVector<mlir::Attribute> deviceTypes;
+ std::optional<mlir::ArrayAttr> existingDeviceTypes =
+ operation.getDeviceTypes();
+
+ // Ensure we keep the existing ones, and in the correct 'new' order.
+ if (existingDeviceTypes) {
+ for (const mlir::Attribute &Attr : *existingDeviceTypes)
+ deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(),
+ cast<mlir::acc::DeviceTypeAttr>(Attr).getValue()));
+ }
- switch (dirKind) {
- case OpenACCDirectiveKind::Init:
- case OpenACCDirectiveKind::Set:
- 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:
+ for (const DeviceTypeArgument &arg : clause.getArchitectures()) {
+ deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(), decodeDeviceType(arg.first)));
+ }
+ operation.removeDeviceTypesAttr();
+ operation.setDeviceTypesAttr(
+ mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
+ } else if constexpr (isOneOfTypes<OpTy, SetOp>) {
+ assert(!operation.getDeviceTypeAttr() && "already have device-type?");
+ assert(clause.getArchitectures().size() <= 1);
+
+ if (!clause.getArchitectures().empty())
+ operation.setDeviceType(
+ decodeDeviceType(clause.getArchitectures()[0].first));
+ } else {
return clauseNotImplemented(clause);
}
}
- // Apply any of the clauses that resulted in an 'attribute'.
- 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.
- // 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
- 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.
- // 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)
- deviceTypes.push_back(
- mlir::acc::DeviceTypeAttr::get(builder.getContext(), DT));
-
- op.setDeviceTypesAttr(
- mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
- } else if constexpr (isOneOfTypes<Op, SetOp>) {
- assert(attrData.deviceTypeArchs.size() <= 1 &&
- "Set can only have a single architecture");
- if (!attrData.deviceTypeArchs.empty())
- op.setDeviceType(attrData.deviceTypeArchs[0]);
+ void VisitSelfClause(const OpenACCSelfClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
+ if (clause.isEmptySelfClause()) {
+ operation.setSelfAttr(true);
+ } else if (clause.isConditionExprClause()) {
+ assert(clause.hasConditionExpr());
+ mlir::Value condition =
+ cgf.evaluateExprAsBool(clause.getConditionExpr());
+
+ mlir::Location exprLoc = cgf.getCIRGenModule().getLoc(
+ clause.getConditionExpr()->getBeginLoc());
+ mlir::IntegerType targetType = mlir::IntegerType::get(
+ &cgf.getMLIRContext(), /*width=*/1,
+ mlir::IntegerType::SignednessSemantics::Signless);
+ auto ConversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
+ exprLoc, targetType, condition);
+ operation.getSelfCondMutable().append(ConversionOp.getResult(0));
} else {
- cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ",
- dirKind);
+ llvm_unreachable("var-list version of self shouldn't get here");
}
+ } else {
+ return clauseNotImplemented(clause);
}
}
};
+template <typename OpTy>
+auto makeClauseEmitter(OpTy &op, CIRGenFunction &cgf, CIRGenBuilderTy &builder,
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc) {
+ return OpenACCClauseCIREmitter<OpTy>(op, cgf, builder, dirKind, dirLoc);
+}
+
} // namespace
template <typename Op, typename TermOp>
@@ -158,24 +170,27 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
-
- // Clause-emitter must be here because it might modify operands.
- 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(builder, op);
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ // Sets insertion point before the 'op', since every new expression needs to
+ // be before the operation.
+ builder.setInsertionPoint(op);
+ makeClauseEmitter(op, *this, builder, dirKind, dirLoc)
+ .VisitClauseList(clauses);
+ }
- mlir::Block &block = op.getRegion().emplaceBlock();
- mlir::OpBuilder::InsertionGuard guardCase(builder);
- builder.setInsertionPointToEnd(&block);
+ {
+ mlir::Block &block = op.getRegion().emplaceBlock();
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ builder.setInsertionPointToEnd(&block);
- LexicalScope ls{*this, start, builder.getInsertionBlock()};
- res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
+ LexicalScope ls{*this, start, builder.getInsertionBlock()};
+ res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
- builder.create<TermOp>(end);
+ builder.create<TermOp>(end);
+ }
return res;
}
@@ -187,14 +202,16 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOp(
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
-
- // Clause-emitter must be here because it might modify operands.
- 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(builder, op);
+
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ // Sets insertion point before the 'op', since every new expression needs to
+ // be before the operation.
+ builder.setInsertionPoint(op);
+ makeClauseEmitter(op, *this, builder, dirKind, dirLoc)
+ .VisitClauseList(clauses);
+ }
return res;
}
diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c b/clang/test/CIR/CodeGenOpenACC/kernels.c
index 0c950fe3d0f9c..934daf9e8ecc0 100644
--- a/clang/test/CIR/CodeGenOpenACC/kernels.c
+++ b/clang/test/CIR/CodeGenOpenACC/kernels.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_kernels(void) {
- // CHECK: cir.func @acc_kernels() {
+void acc_kernels(int cond) {
+ // CHECK: cir.func @acc_kernels(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc kernels
{}
@@ -38,5 +40,29 @@ void acc_kernels(void) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT:}
+#pragma acc kernels self
+ {}
+ // CHECK-NEXT: acc.kernels {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc kernels self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
+#pragma acc kernels self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c b/clang/test/CIR/CodeGenOpenACC/parallel.c
index e18270435460c..c7a4bda6faa74 100644
--- a/clang/test/CIR/CodeGenOpenACC/parallel.c
+++ b/clang/test/CIR/CodeGenOpenACC/parallel.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_parallel(void) {
- // CHECK: cir.func @acc_parallel() {
+void acc_parallel(int cond) {
+ // CHECK: cir.func @acc_parallel(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc parallel
{}
// CHECK-NEXT: acc.parallel {
@@ -37,5 +39,29 @@ void acc_parallel(void) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc parallel self
+ {}
+ // CHECK-NEXT: acc.parallel {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc parallel self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
+#pragma acc parallel self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c b/clang/test/CIR/CodeGenOpenACC/serial.c
index 72a0995549da3..38a38ad6c9514 100644
--- a/clang/test/CIR/CodeGenOpenACC/serial.c
+++ b/clang/test/CIR/CodeGenOpenACC/serial.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_serial(void) {
- // CHECK: cir.func @acc_serial() {
+void acc_serial(int cond) {
+ // CHECK: cir.func @acc_serial(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc serial
{}
@@ -38,5 +40,29 @@ void acc_serial(void) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc serial self
+ {}
+ // CHECK-NEXT: acc.serial {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc serial self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
+#pragma acc serial self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
|
|
@llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) ChangesThis is our first attempt at lowering a clause that is an 'operand' in the OpenACC operand, so it does quite a bit of refactoring. My previous plans on how to emit the clauses was not viable, so we instead do 'create the op, then use the visitor to fill in the operands'. This resulted in the 'applyAttributes' function getting removed and a few other functions simplified. Additionally, it requires setting the insertion point a little to make sure we're inserting 'around' the operation correctly. Finally, since the OpenACC dialect only understands the MLIR types, we had to introduce a use of the unrealized-conversion-cast, which we'll probably getting good use out of in the future. Full diff: https://github.com/llvm/llvm-project/pull/135851.diff 5 Files Affected:
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 3687af76a559f..681567228cbb0 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -430,6 +430,11 @@ class OpenACCSelfClause final
}
bool isConditionExprClause() const { return HasConditionExpr.has_value(); }
+ bool isVarListClause() const { return !isConditionExprClause(); }
+ bool isEmptySelfClause() const {
+ return (isConditionExprClause() && !hasConditionExpr()) ||
+ (!isConditionExprClause() && getVarList().empty());
+ }
bool hasConditionExpr() const {
assert(HasConditionExpr.has_value() &&
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 152f996ed0fed..92f3ad2a68eb8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -32,46 +32,52 @@ constexpr bool isOneOfTypes =
template <typename ToTest, typename T>
constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
+template <typename OpTy>
class OpenACCClauseCIREmitter final
- : public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
- CIRGenModule &cgm;
+ : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> {
+ OpTy &operation;
+ CIRGenFunction &cgf;
+ CIRGenBuilderTy &builder;
+
// This is necessary since a few of the clauses emit differently based on the
// directive kind they are attached to.
OpenACCDirectiveKind dirKind;
+ // This source location should be able to go away once the NYI diagnostics are
+ // gone.
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) {
- cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
+ cgf.getCIRGenModule().errorNYI(c.getSourceRange(), "OpenACC Clause",
+ c.getClauseKind());
}
public:
- OpenACCClauseCIREmitter(CIRGenModule &cgm, OpenACCDirectiveKind dirKind,
- SourceLocation dirLoc)
- : cgm(cgm), dirKind(dirKind), dirLoc(dirLoc) {}
+ OpenACCClauseCIREmitter(OpTy &operation, CIRGenFunction &cgf,
+ CIRGenBuilderTy &builder,
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc)
+ : operation(operation), cgf(cgf), builder(builder), dirKind(dirKind),
+ dirLoc(dirLoc) {}
void VisitClause(const OpenACCClause &clause) {
clauseNotImplemented(clause);
}
void VisitDefaultClause(const OpenACCDefaultClause &clause) {
- switch (clause.getDefaultClauseKind()) {
- case OpenACCDefaultClauseKind::None:
- attrData.defaultVal = ClauseDefaultValue::None;
- break;
- case OpenACCDefaultClauseKind::Present:
- attrData.defaultVal = ClauseDefaultValue::Present;
- break;
- case OpenACCDefaultClauseKind::Invalid:
- break;
+ // 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<OpTy, ParallelOp, SerialOp, KernelsOp, DataOp>) {
+ switch (clause.getDefaultClauseKind()) {
+ case OpenACCDefaultClauseKind::None:
+ operation.setDefaultAttr(ClauseDefaultValue::None);
+ break;
+ case OpenACCDefaultClauseKind::Present:
+ operation.setDefaultAttr(ClauseDefaultValue::Present);
+ break;
+ case OpenACCDefaultClauseKind::Invalid:
+ break;
+ }
+ } else {
+ return clauseNotImplemented(clause);
}
}
@@ -89,64 +95,70 @@ class OpenACCClauseCIREmitter final
}
void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) {
+ llvm::SmallVector<mlir::Attribute> deviceTypes;
+ std::optional<mlir::ArrayAttr> existingDeviceTypes =
+ operation.getDeviceTypes();
+
+ // Ensure we keep the existing ones, and in the correct 'new' order.
+ if (existingDeviceTypes) {
+ for (const mlir::Attribute &Attr : *existingDeviceTypes)
+ deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(),
+ cast<mlir::acc::DeviceTypeAttr>(Attr).getValue()));
+ }
- switch (dirKind) {
- case OpenACCDirectiveKind::Init:
- case OpenACCDirectiveKind::Set:
- 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:
+ for (const DeviceTypeArgument &arg : clause.getArchitectures()) {
+ deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(), decodeDeviceType(arg.first)));
+ }
+ operation.removeDeviceTypesAttr();
+ operation.setDeviceTypesAttr(
+ mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
+ } else if constexpr (isOneOfTypes<OpTy, SetOp>) {
+ assert(!operation.getDeviceTypeAttr() && "already have device-type?");
+ assert(clause.getArchitectures().size() <= 1);
+
+ if (!clause.getArchitectures().empty())
+ operation.setDeviceType(
+ decodeDeviceType(clause.getArchitectures()[0].first));
+ } else {
return clauseNotImplemented(clause);
}
}
- // Apply any of the clauses that resulted in an 'attribute'.
- 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.
- // 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
- 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.
- // 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)
- deviceTypes.push_back(
- mlir::acc::DeviceTypeAttr::get(builder.getContext(), DT));
-
- op.setDeviceTypesAttr(
- mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
- } else if constexpr (isOneOfTypes<Op, SetOp>) {
- assert(attrData.deviceTypeArchs.size() <= 1 &&
- "Set can only have a single architecture");
- if (!attrData.deviceTypeArchs.empty())
- op.setDeviceType(attrData.deviceTypeArchs[0]);
+ void VisitSelfClause(const OpenACCSelfClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
+ if (clause.isEmptySelfClause()) {
+ operation.setSelfAttr(true);
+ } else if (clause.isConditionExprClause()) {
+ assert(clause.hasConditionExpr());
+ mlir::Value condition =
+ cgf.evaluateExprAsBool(clause.getConditionExpr());
+
+ mlir::Location exprLoc = cgf.getCIRGenModule().getLoc(
+ clause.getConditionExpr()->getBeginLoc());
+ mlir::IntegerType targetType = mlir::IntegerType::get(
+ &cgf.getMLIRContext(), /*width=*/1,
+ mlir::IntegerType::SignednessSemantics::Signless);
+ auto ConversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
+ exprLoc, targetType, condition);
+ operation.getSelfCondMutable().append(ConversionOp.getResult(0));
} else {
- cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ",
- dirKind);
+ llvm_unreachable("var-list version of self shouldn't get here");
}
+ } else {
+ return clauseNotImplemented(clause);
}
}
};
+template <typename OpTy>
+auto makeClauseEmitter(OpTy &op, CIRGenFunction &cgf, CIRGenBuilderTy &builder,
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc) {
+ return OpenACCClauseCIREmitter<OpTy>(op, cgf, builder, dirKind, dirLoc);
+}
+
} // namespace
template <typename Op, typename TermOp>
@@ -158,24 +170,27 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
-
- // Clause-emitter must be here because it might modify operands.
- 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(builder, op);
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ // Sets insertion point before the 'op', since every new expression needs to
+ // be before the operation.
+ builder.setInsertionPoint(op);
+ makeClauseEmitter(op, *this, builder, dirKind, dirLoc)
+ .VisitClauseList(clauses);
+ }
- mlir::Block &block = op.getRegion().emplaceBlock();
- mlir::OpBuilder::InsertionGuard guardCase(builder);
- builder.setInsertionPointToEnd(&block);
+ {
+ mlir::Block &block = op.getRegion().emplaceBlock();
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ builder.setInsertionPointToEnd(&block);
- LexicalScope ls{*this, start, builder.getInsertionBlock()};
- res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
+ LexicalScope ls{*this, start, builder.getInsertionBlock()};
+ res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
- builder.create<TermOp>(end);
+ builder.create<TermOp>(end);
+ }
return res;
}
@@ -187,14 +202,16 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOp(
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
-
- // Clause-emitter must be here because it might modify operands.
- 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(builder, op);
+
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ // Sets insertion point before the 'op', since every new expression needs to
+ // be before the operation.
+ builder.setInsertionPoint(op);
+ makeClauseEmitter(op, *this, builder, dirKind, dirLoc)
+ .VisitClauseList(clauses);
+ }
return res;
}
diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c b/clang/test/CIR/CodeGenOpenACC/kernels.c
index 0c950fe3d0f9c..934daf9e8ecc0 100644
--- a/clang/test/CIR/CodeGenOpenACC/kernels.c
+++ b/clang/test/CIR/CodeGenOpenACC/kernels.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_kernels(void) {
- // CHECK: cir.func @acc_kernels() {
+void acc_kernels(int cond) {
+ // CHECK: cir.func @acc_kernels(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc kernels
{}
@@ -38,5 +40,29 @@ void acc_kernels(void) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT:}
+#pragma acc kernels self
+ {}
+ // CHECK-NEXT: acc.kernels {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc kernels self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
+#pragma acc kernels self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c b/clang/test/CIR/CodeGenOpenACC/parallel.c
index e18270435460c..c7a4bda6faa74 100644
--- a/clang/test/CIR/CodeGenOpenACC/parallel.c
+++ b/clang/test/CIR/CodeGenOpenACC/parallel.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_parallel(void) {
- // CHECK: cir.func @acc_parallel() {
+void acc_parallel(int cond) {
+ // CHECK: cir.func @acc_parallel(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc parallel
{}
// CHECK-NEXT: acc.parallel {
@@ -37,5 +39,29 @@ void acc_parallel(void) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc parallel self
+ {}
+ // CHECK-NEXT: acc.parallel {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc parallel self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
+#pragma acc parallel self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c b/clang/test/CIR/CodeGenOpenACC/serial.c
index 72a0995549da3..38a38ad6c9514 100644
--- a/clang/test/CIR/CodeGenOpenACC/serial.c
+++ b/clang/test/CIR/CodeGenOpenACC/serial.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_serial(void) {
- // CHECK: cir.func @acc_serial() {
+void acc_serial(int cond) {
+ // CHECK: cir.func @acc_serial(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc serial
{}
@@ -38,5 +40,29 @@ void acc_serial(void) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc serial self
+ {}
+ // CHECK-NEXT: acc.serial {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc serial self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
+#pragma acc serial self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
|
erichkeane
commented
Apr 15, 2025
| bool isVarListClause() const { return !isConditionExprClause(); } | ||
| bool isEmptySelfClause() const { | ||
| return (isConditionExprClause() && !hasConditionExpr()) || | ||
| (!isConditionExprClause() && getVarList().empty()); |
There was a problem hiding this comment.
These convenience functions made a lot of sense/got good use in codegen, so they seemed sensible to add.
| @@ -32,46 +32,52 @@ constexpr bool isOneOfTypes = | |||
| template <typename ToTest, typename T> | |||
| constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>; | |||
|
|
|||
| template <typename OpTy> | |||
There was a problem hiding this comment.
We end up needing to template-ize this, since the visitor handler functions are not needing to if-constexpr instead of just the apply function.
| CIRGenModule &cgm; | ||
| : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> { | ||
| OpTy &operation; | ||
| CIRGenFunction &cgf; |
There was a problem hiding this comment.
Function instead of Module because we need to be able to emit expressions, which we do at the function.
| mlir::IntegerType targetType = mlir::IntegerType::get( | ||
| &cgf.getMLIRContext(), /*width=*/1, | ||
| mlir::IntegerType::SignednessSemantics::Signless); | ||
| auto ConversionOp = builder.create<mlir::UnrealizedConversionCastOp>( |
There was a problem hiding this comment.
This is the magic to convert a cir.bool type to a mlir::IntegerType, which @bcardosolopes says is handled during one of our lowering steps.
| } | ||
| } | ||
| }; | ||
|
|
||
| template <typename OpTy> | ||
| auto makeClauseEmitter(OpTy &op, CIRGenFunction &cgf, CIRGenBuilderTy &builder, |
There was a problem hiding this comment.
We need this because we have a few host-compilers that don't support CTAD fully yet.
| mlir::OpBuilder::InsertionGuard guardCase(builder); | ||
| // Sets insertion point before the 'op', since every new expression needs to | ||
| // be before the operation. | ||
| builder.setInsertionPoint(op); |
There was a problem hiding this comment.
See here for the magic around getting the 'insertion' point right. We now only view the clauses information 1x in the emitter (rather than collect data, then emit data for attributes), but all of the ops they 'insert' need to happen before the OpenACC operation.
|
+Razvan/Valentin for review of the OpenACC-IR. |
clementval
reviewed
Apr 15, 2025
andykaylor
approved these changes
Apr 15, 2025
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/190/builds/18368 Here is the relevant piece of the build log for the reference |
var-const
pushed a commit
to ldionne/llvm-project
that referenced
this pull request
Apr 17, 2025
…135851) This is our first attempt at lowering a clause that is an 'operand' in the OpenACC operand, so it does quite a bit of refactoring. My previous plans on how to emit the clauses was not viable, so we instead do 'create the op, then use the visitor to fill in the operands'. This resulted in the 'applyAttributes' function getting removed and a few other functions simplified. Additionally, it requires setting the insertion point a little to make sure we're inserting 'around' the operation correctly. Finally, since the OpenACC dialect only understands the MLIR types, we had to introduce a use of the unrealized-conversion-cast, which we'll probably getting good use out of in the future.
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.
This is our first attempt at lowering a clause that is an 'operand' in the OpenACC operand, so it does quite a bit of refactoring. My previous plans on how to emit the clauses was not viable, so we instead do 'create the op, then use the visitor to fill in the operands'. This resulted in the 'applyAttributes' function getting removed and a few other functions simplified.
Additionally, it requires setting the insertion point a little to make sure we're inserting 'around' the operation correctly.
Finally, since the OpenACC dialect only understands the MLIR types, we had to introduce a use of the unrealized-conversion-cast, which we'll probably getting good use out of in the future.