[clang] fa273e1 - [OpenACC][CIR] Implement 'data' construct lowering (#135038)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 9 10:45:20 PDT 2025
Author: Erich Keane
Date: 2025-04-09T10:45:17-07:00
New Revision: fa273e1158edb109e1c392a0d8e18b711d0e008e
URL: https://github.com/llvm/llvm-project/commit/fa273e1158edb109e1c392a0d8e18b711d0e008e
DIFF: https://github.com/llvm/llvm-project/commit/fa273e1158edb109e1c392a0d8e18b711d0e008e.diff
LOG: [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.
Added:
clang/test/CIR/CodeGenOpenACC/data.c
Modified:
clang/include/clang/AST/OpenACCClause.h
clang/lib/CIR/CodeGen/CIRGenFunction.h
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
clang/test/CIR/CodeGenOpenACC/kernels.c
clang/test/CIR/CodeGenOpenACC/parallel.c
clang/test/CIR/CodeGenOpenACC/serial.c
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index fda1837594c99..3687af76a559f 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -1316,11 +1316,13 @@ template <class Impl> class OpenACCClauseVisitor {
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 @@ template <class Impl> class OpenACCClauseVisitor {
#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 fb5ec6a868a1b..c30fcc2a05f87 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -571,14 +571,13 @@ class CIRGenFunction : public CIRGenTypeCache {
// 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 7a8879add784a..e7e56d3602e3a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -27,6 +27,12 @@ class OpenACCClauseCIREmitter final
: 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 @@ class OpenACCClauseCIREmitter final
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);
+ }
+
+ 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);
}
-#include "clang/Basic/OpenACCClauses.def"
};
} // 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();
+ 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);
- llvm::SmallVector<mlir::Type> retTy;
- llvm::SmallVector<mlir::Value> operands;
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,19 +101,28 @@ CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) {
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");
}
}
+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");
@@ -97,11 +133,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
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 0000000000000..9e636f68aad1b
--- /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 91684859f7115..0c950fe3d0f9c 100644
--- a/clang/test/CIR/CodeGenOpenACC/kernels.c
+++ b/clang/test/CIR/CodeGenOpenACC/kernels.c
@@ -6,9 +6,21 @@ void acc_kernels(void) {
{}
// 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 @@ void acc_kernels(void) {
// 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 7c1509a129980..e18270435460c 100644
--- a/clang/test/CIR/CodeGenOpenACC/parallel.c
+++ b/clang/test/CIR/CodeGenOpenACC/parallel.c
@@ -5,9 +5,21 @@ void acc_parallel(void) {
#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 @@ void acc_parallel(void) {
// 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 9897cd3d4e8d9..72a0995549da3 100644
--- a/clang/test/CIR/CodeGenOpenACC/serial.c
+++ b/clang/test/CIR/CodeGenOpenACC/serial.c
@@ -6,9 +6,21 @@ void acc_serial(void) {
{}
// 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 @@ void acc_serial(void) {
// CHECK-NEXT: }
// cir.scope end:
// CHECK-NEXT: }
- // CHECK-NEXT:acc.yield
+ // CHECK-NEXT: acc.yield
// CHECK-NEXT:}
// CHECK-NEXT: cir.return
More information about the cfe-commits
mailing list