[OpenACC][CIR] Implement 'data' construct lowering (#135038)
This patch does the lowering of the OpenACC 'data' construct, which
requires getting the `default` clause (as `data` requires at least 1 of
a list of clauses, and this is the easiest one). The lowering of the
clauses appears to happen in 1 of 2 ways: a- as an operand. or b- as an
attribute.
This patch adds infrastructure to lower as an attribute, as that is how
'data' works.
In addition to that, it changes the OpenACCClauseVisitor a bit, which
previously just required that each of the derived classes have all of
the clauses covered. This patch modifies it so that the visitor directly
calls the derived class from its visitor function, which leaves the
base-class ones the ability to defer to a generic function. This was
previously like this because I had some use cases that I didn't end up
using, and the 'generic' function here seems much more useful.
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index fda1837..3687af7 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -1316,11 +1316,13 @@
switch (C->getClauseKind()) {
#define VISIT_CLAUSE(CLAUSE_NAME) \
case OpenACCClauseKind::CLAUSE_NAME: \
- Visit##CLAUSE_NAME##Clause(*cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
+ getDerived().Visit##CLAUSE_NAME##Clause( \
+ *cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
return;
#define CLAUSE_ALIAS(ALIAS_NAME, CLAUSE_NAME, DEPRECATED) \
case OpenACCClauseKind::ALIAS_NAME: \
- Visit##CLAUSE_NAME##Clause(*cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
+ getDerived().Visit##CLAUSE_NAME##Clause( \
+ *cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
return;
#include "clang/Basic/OpenACCClauses.def"
@@ -1333,7 +1335,7 @@
#define VISIT_CLAUSE(CLAUSE_NAME) \
void Visit##CLAUSE_NAME##Clause( \
const OpenACC##CLAUSE_NAME##Clause &Clause) { \
- return getDerived().Visit##CLAUSE_NAME##Clause(Clause); \
+ return getDerived().VisitClause(Clause); \
}
#include "clang/Basic/OpenACCClauses.def"
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index fb5ec6a..c30fcc2 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -571,14 +571,13 @@
// OpenACC Emission
//===--------------------------------------------------------------------===//
private:
- // Function to do the basic implementation of a 'compute' operation, including
- // the clauses/etc. This might be generalizable in the future to work for
- // other constructs, or at least be the base for construct emission.
+ // Function to do the basic implementation of an operation with an Associated
+ // Statement. Models AssociatedStmtConstruct.
template <typename Op, typename TermOp>
mlir::LogicalResult
- emitOpenACCComputeOp(mlir::Location start, mlir::Location end,
- llvm::ArrayRef<const OpenACCClause *> clauses,
- const Stmt *structuredBlock);
+ emitOpenACCOpAssociatedStmt(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 7a8879a..e7e56d3 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -27,6 +27,12 @@
: public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
CIRGenModule &cgm;
+ struct AttributeData {
+ // Value of the 'default' attribute, added on 'data' and 'compute'/etc
+ // constructs as a 'default-attr'.
+ std::optional<ClauseDefaultValue> defaultVal = std::nullopt;
+ } attrData;
+
void clauseNotImplemented(const OpenACCClause &c) {
cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
}
@@ -34,34 +40,55 @@
public:
OpenACCClauseCIREmitter(CIRGenModule &cgm) : cgm(cgm) {}
-#define VISIT_CLAUSE(CN) \
- void Visit##CN##Clause(const OpenACC##CN##Clause &clause) { \
- clauseNotImplemented(clause); \
+ void VisitClause(const OpenACCClause &clause) {
+ clauseNotImplemented(clause);
}
-#include "clang/Basic/OpenACCClauses.def"
+
+ 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;
+ }
+ }
+
+ // 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);
+ }
};
} // namespace
template <typename Op, typename TermOp>
-mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp(
+mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
mlir::Location start, mlir::Location end,
- llvm::ArrayRef<const OpenACCClause *> clauses,
- const Stmt *structuredBlock) {
+ llvm::ArrayRef<const OpenACCClause *> clauses, const Stmt *associatedStmt) {
mlir::LogicalResult res = mlir::success();
- OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
- clauseEmitter.VisitClauseList(clauses);
-
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
+
+ // Clause-emitter must be here because it might modify operands.
+ OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
+ clauseEmitter.VisitClauseList(clauses);
+
auto op = builder.create<Op>(start, retTy, operands);
+ // Apply the attributes derived from the clauses.
+ clauseEmitter.applyAttributes(op);
+
mlir::Block &block = op.getRegion().emplaceBlock();
mlir::OpBuilder::InsertionGuard guardCase(builder);
builder.setInsertionPointToEnd(&block);
LexicalScope ls{*this, start, builder.getInsertionBlock()};
- res = emitStmt(structuredBlock, /*useCurrentScope=*/true);
+ res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
builder.create<TermOp>(end);
return res;
@@ -74,13 +101,13 @@
switch (s.getDirectiveKind()) {
case OpenACCDirectiveKind::Parallel:
- return emitOpenACCComputeOp<ParallelOp, mlir::acc::YieldOp>(
+ return emitOpenACCOpAssociatedStmt<ParallelOp, mlir::acc::YieldOp>(
start, end, s.clauses(), s.getStructuredBlock());
case OpenACCDirectiveKind::Serial:
- return emitOpenACCComputeOp<SerialOp, mlir::acc::YieldOp>(
+ return emitOpenACCOpAssociatedStmt<SerialOp, mlir::acc::YieldOp>(
start, end, s.clauses(), s.getStructuredBlock());
case OpenACCDirectiveKind::Kernels:
- return emitOpenACCComputeOp<KernelsOp, mlir::acc::TerminatorOp>(
+ return emitOpenACCOpAssociatedStmt<KernelsOp, mlir::acc::TerminatorOp>(
start, end, s.clauses(), s.getStructuredBlock());
default:
llvm_unreachable("invalid compute construct kind");
@@ -88,6 +115,15 @@
}
mlir::LogicalResult
+CIRGenFunction::emitOpenACCDataConstruct(const OpenACCDataConstruct &s) {
+ mlir::Location start = getLoc(s.getSourceRange().getEnd());
+ mlir::Location end = getLoc(s.getSourceRange().getEnd());
+
+ return emitOpenACCOpAssociatedStmt<DataOp, mlir::acc::TerminatorOp>(
+ start, end, s.clauses(), s.getStructuredBlock());
+}
+
+mlir::LogicalResult
CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) {
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Loop Construct");
return mlir::failure();
@@ -97,11 +133,6 @@
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Combined Construct");
return mlir::failure();
}
-mlir::LogicalResult
-CIRGenFunction::emitOpenACCDataConstruct(const OpenACCDataConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Data Construct");
- return mlir::failure();
-}
mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
const OpenACCEnterDataConstruct &s) {
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
diff --git a/clang/test/CIR/CodeGenOpenACC/data.c b/clang/test/CIR/CodeGenOpenACC/data.c
new file mode 100644
index 0000000..9e636f6
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/data.c
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_data(void) {
+ // CHECK: cir.func @acc_data() {
+
+#pragma acc data default(none)
+ {
+ int i = 0;
+ ++i;
+ }
+ // CHECK-NEXT: acc.data {
+ // CHECK-NEXT: cir.alloca
+ // CHECK-NEXT: cir.const
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ // CHECK-NEXT: cir.unary
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(present)
+ {
+ int i = 0;
+ ++i;
+ }
+ // CHECK-NEXT: acc.data {
+ // CHECK-NEXT: cir.alloca
+ // CHECK-NEXT: cir.const
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ // CHECK-NEXT: cir.unary
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}
+
+ // CHECK-NEXT: cir.return
+}
diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c b/clang/test/CIR/CodeGenOpenACC/kernels.c
index 9168485..0c950fe 100644
--- a/clang/test/CIR/CodeGenOpenACC/kernels.c
+++ b/clang/test/CIR/CodeGenOpenACC/kernels.c
@@ -6,9 +6,21 @@
{}
// CHECK-NEXT: acc.kernels {
- // CHECK-NEXT:acc.terminator
+ // CHECK-NEXT: acc.terminator
// CHECK-NEXT:}
+#pragma acc kernels default(none)
+ {}
+ // CHECK-NEXT: acc.kernels {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc kernels default(present)
+ {}
+ // CHECK-NEXT: acc.kernels {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}
+
#pragma acc kernels
while(1){}
// CHECK-NEXT: acc.kernels {
@@ -23,7 +35,7 @@
// CHECK-NEXT: }
// cir.scope end:
// CHECK-NEXT: }
- // CHECK-NEXT:acc.terminator
+ // CHECK-NEXT: acc.terminator
// CHECK-NEXT:}
// CHECK-NEXT: cir.return
diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c b/clang/test/CIR/CodeGenOpenACC/parallel.c
index 7c1509a..e182704354 100644
--- a/clang/test/CIR/CodeGenOpenACC/parallel.c
+++ b/clang/test/CIR/CodeGenOpenACC/parallel.c
@@ -5,9 +5,21 @@
#pragma acc parallel
{}
// CHECK-NEXT: acc.parallel {
- // CHECK-NEXT:acc.yield
+ // CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc parallel default(none)
+ {}
+ // CHECK-NEXT: acc.parallel {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc parallel default(present)
+ {}
+ // CHECK-NEXT: acc.parallel {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}
+
#pragma acc parallel
while(1){}
// CHECK-NEXT: acc.parallel {
@@ -22,7 +34,7 @@
// CHECK-NEXT: }
// cir.scope end:
// CHECK-NEXT: }
- // CHECK-NEXT:acc.yield
+ // CHECK-NEXT: acc.yield
// CHECK-NEXT:}
// CHECK-NEXT: cir.return
diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c b/clang/test/CIR/CodeGenOpenACC/serial.c
index 9897cd3..72a0995 100644
--- a/clang/test/CIR/CodeGenOpenACC/serial.c
+++ b/clang/test/CIR/CodeGenOpenACC/serial.c
@@ -6,9 +6,21 @@
{}
// CHECK-NEXT: acc.serial {
- // CHECK-NEXT:acc.yield
+ // CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc serial default(none)
+ {}
+ // CHECK-NEXT: acc.serial {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc serial default(present)
+ {}
+ // CHECK-NEXT: acc.serial {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}
+
#pragma acc serial
while(1){}
// CHECK-NEXT: acc.serial {
@@ -23,7 +35,7 @@
// CHECK-NEXT: }
// cir.scope end:
// CHECK-NEXT: }
- // CHECK-NEXT:acc.yield
+ // CHECK-NEXT: acc.yield
// CHECK-NEXT:}
// CHECK-NEXT: cir.return
