[OpenACC][CIR] Initial patch to do OpenACC->IR lowering #134936
Merged
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 patch adds some lowering code for Compute Constructs, plus the infrastructure to someday do clauses. Doing this requires adding the dialect to the CIRGenerator. This patch does not however implement/correctly initialize lowering from OpenACC-Dialect to anything lower however.
llvmbot
added
clang
|
@llvm/pr-subscribers-clang @llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) ChangesThis patch adds some lowering code for Compute Constructs, plus the infrastructure to someday do clauses. Doing this requires adding the dialect to the CIRGenerator. This patch does not however implement/correctly initialize lowering from OpenACC-Dialect to anything lower however. Full diff: https://github.com/llvm/llvm-project/pull/134936.diff 8 Files Affected:
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index b3a5746af7cb0..fda1837594c99 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -38,6 +38,7 @@ class OpenACCClause {
OpenACCClauseKind getClauseKind() const { return Kind; }
SourceLocation getBeginLoc() const { return Location.getBegin(); }
SourceLocation getEndLoc() const { return Location.getEnd(); }
+ SourceRange getSourceRange() const { return Location; }
static bool classof(const OpenACCClause *) { return true; }
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index dde665a75ae57..3628f0ea4510e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -536,6 +536,16 @@ 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.
+ template <typename Op, typename Terminator>
+ mlir::LogicalResult
+ emitOpenACCComputeOp(mlir::Location start,
+ mlir::Location end,
+ llvm::ArrayRef<const OpenACCClause *> clauses,
+ const Stmt *structuredBlock);
public:
mlir::LogicalResult
emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s);
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index cbae170162ffe..cccf7ee826ea6 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -12,16 +12,84 @@
#include "CIRGenBuilder.h"
#include "CIRGenFunction.h"
+#include "clang/AST/OpenACCClause.h"
#include "clang/AST/StmtOpenACC.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+
using namespace clang;
using namespace clang::CIRGen;
using namespace cir;
+using namespace mlir::acc;
+
+namespace {
+class OpenACCClauseCIREmitter final
+ : public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
+ CIRGenModule &cgm;
+
+ void clauseNotImplemented(const OpenACCClause &c) {
+ cgm.errorNYI(c.getSourceRange(), "OpenACC Clause",
+ c.getClauseKind());
+ }
+
+ public:
+ OpenACCClauseCIREmitter(CIRGenModule &cgm): cgm(cgm){}
+
+ void VisitClauseList(llvm::ArrayRef<const OpenACCClause *> clauses) {
+ for (auto *clause : clauses)
+ Visit(clause);
+ }
+
+#define VISIT_CLAUSE(CN) \
+ void Visit##CN##Clause(const OpenACC##CN##Clause &clause) { \
+ clauseNotImplemented(clause); \
+ }
+#include "clang/Basic/OpenACCClauses.def"
+ };
+}
+
+template <typename Op, typename Terminator>
+mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp(
+ mlir::Location start, mlir::Location end,
+ llvm::ArrayRef<const OpenACCClause *> clauses,
+ const Stmt *structuredBlock) {
+ mlir::LogicalResult res = mlir::success();
+
+ OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
+
+ llvm::SmallVector<mlir::Type> retTy;
+ llvm::SmallVector<mlir::Value> operands;
+ auto op = builder.create<Op>(start, retTy, operands);
+
+ 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);
+
+ builder.create<Terminator>(end);
+ return res;
+}
mlir::LogicalResult
CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Compute Construct");
- return mlir::failure();
+ auto start = getLoc(s.getSourceRange().getEnd());
+ auto end = getLoc(s.getSourceRange().getEnd());
+
+ switch (s.getDirectiveKind()) {
+ case OpenACCDirectiveKind::Parallel:
+ return emitOpenACCComputeOp<ParallelOp, mlir::acc::YieldOp>(
+ start, end, s.clauses(), s.getStructuredBlock());
+ case OpenACCDirectiveKind::Serial:
+ return emitOpenACCComputeOp<SerialOp, mlir::acc::YieldOp>(
+ start, end, s.clauses(), s.getStructuredBlock());
+ case OpenACCDirectiveKind::Kernels:
+ return emitOpenACCComputeOp<KernelsOp, mlir::acc::TerminatorOp>(
+ start, end, s.clauses(), s.getStructuredBlock());
+ default:
+ llvm_unreachable("invalid compute construct kind");
+ }
}
mlir::LogicalResult
diff --git a/clang/lib/CIR/CodeGen/CIRGenerator.cpp b/clang/lib/CIR/CodeGen/CIRGenerator.cpp
index 33f0c292c7710..aa3864deb733c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenerator.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenerator.cpp
@@ -12,6 +12,7 @@
#include "CIRGenModule.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/MLIRContext.h"
#include "clang/AST/DeclGroup.h"
@@ -36,6 +37,7 @@ void CIRGenerator::Initialize(ASTContext &astContext) {
mlirContext = std::make_unique<mlir::MLIRContext>();
mlirContext->loadDialect<cir::CIRDialect>();
+ mlirContext->getOrLoadDialect<mlir::acc::OpenACCDialect>();
cgm = std::make_unique<clang::CIRGen::CIRGenModule>(
*mlirContext.get(), astContext, codeGenOpts, diags);
}
diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c b/clang/test/CIR/CodeGenOpenACC/kernels.c
new file mode 100644
index 0000000000000..91684859f7115
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/kernels.c
@@ -0,0 +1,30 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_kernels(void) {
+ // CHECK: cir.func @acc_kernels() {
+#pragma acc kernels
+ {}
+
+ // CHECK-NEXT: acc.kernels {
+ // CHECK-NEXT:acc.terminator
+ // CHECK-NEXT:}
+
+#pragma acc kernels
+ while(1){}
+ // CHECK-NEXT: acc.kernels {
+ // CHECK-NEXT: cir.scope {
+ // CHECK-NEXT: cir.while {
+ // CHECK-NEXT: %[[INT:.*]] = cir.const #cir.int<1>
+ // CHECK-NEXT: %[[CAST:.*]] = cir.cast(int_to_bool, %[[INT]] :
+ // CHECK-NEXT: cir.condition(%[[CAST]])
+ // CHECK-NEXT: } do {
+ // CHECK-NEXT: cir.yield
+ // cir.while do end:
+ // CHECK-NEXT: }
+ // cir.scope end:
+ // CHECK-NEXT: }
+ // CHECK-NEXT:acc.terminator
+ // CHECK-NEXT:}
+
+ // CHECK-NEXT: cir.return
+}
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
index 61bed79dc14ea..a7a179c0b2e3c 100644
--- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
@@ -3,9 +3,9 @@
void HelloWorld(int *A, int *B, int *C, int N) {
-// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Compute Construct}}
+// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Combined Construct}}
// expected-error@+1{{ClangIR code gen Not Yet Implemented: statement}}
-#pragma acc parallel
+#pragma acc parallel loop
for (unsigned I = 0; I < N; ++I)
A[I] = B[I] + C[I];
diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c b/clang/test/CIR/CodeGenOpenACC/parallel.c
new file mode 100644
index 0000000000000..7c1509a129980
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/parallel.c
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_parallel(void) {
+ // CHECK: cir.func @acc_parallel() {
+#pragma acc parallel
+ {}
+ // CHECK-NEXT: acc.parallel {
+ // CHECK-NEXT:acc.yield
+ // CHECK-NEXT:}
+
+#pragma acc parallel
+ while(1){}
+ // CHECK-NEXT: acc.parallel {
+ // CHECK-NEXT: cir.scope {
+ // CHECK-NEXT: cir.while {
+ // CHECK-NEXT: %[[INT:.*]] = cir.const #cir.int<1>
+ // CHECK-NEXT: %[[CAST:.*]] = cir.cast(int_to_bool, %[[INT]] :
+ // CHECK-NEXT: cir.condition(%[[CAST]])
+ // CHECK-NEXT: } do {
+ // CHECK-NEXT: cir.yield
+ // cir.while do end:
+ // CHECK-NEXT: }
+ // cir.scope end:
+ // CHECK-NEXT: }
+ // 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
new file mode 100644
index 0000000000000..690fc7ad33e63
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/serial.c
@@ -0,0 +1,30 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_serial(void) {
+ // CHECK: cir.func @acc_serial() {
+#pragma acc serial
+ {}
+
+ // CHECK-NEXT: acc.serial {
+ // CHECK-NEXT:acc.yield
+ // CHECK-NEXT:}
+
+#pragma acc serial
+ while(1){}
+ // CHECK-NEXT: acc.serial {
+ // CHECK-NEXT: cir.scope {
+ // CHECK-NEXT: cir.while {
+ // CHECK-NEXT: %[[INT:.*]] = cir.const #cir.int<1>
+ // CHECK-NEXT: %[[CAST:.*]] = cir.cast(int_to_bool, %[[INT]] :
+ // CHECK-NEXT: cir.condition(%[[CAST]])
+ // CHECK-NEXT: } do {
+ // CHECK-NEXT: cir.yield
+ // cir.while do end:
+ // CHECK-NEXT: }
+ // cir.scope end:
+ // CHECK-NEXT: }
+ // CHECK-NEXT:acc.terminator
+ // CHECK-NEXT:}
+
+ // CHECK-NEXT: cir.return
+}
|
andykaylor
approved these changes
Apr 8, 2025
There was a problem hiding this comment.
This looks good except for a few small complaints. It might be good to have someone more familiar with the OpenACC dialect review the tests. @clementval perhaps?
| mlir::OpBuilder::InsertionGuard guardCase(builder); | ||
| builder.setInsertionPointToEnd(&block); | ||
|
|
||
| LexicalScope LS{*this, start, builder.getInsertionBlock()}; |
There was a problem hiding this comment.
Suggested change
| LexicalScope LS{*this, start, builder.getInsertionBlock()}; | |
| LexicalScope ls{*this, start, builder.getInsertionBlock()}; |
|
|
||
| mlir::LogicalResult | ||
| CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) { | ||
| getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Compute Construct"); | ||
| return mlir::failure(); | ||
| auto start = getLoc(s.getSourceRange().getEnd()); |
There was a problem hiding this comment.
Suggested change
| auto start = getLoc(s.getSourceRange().getEnd()); | |
| mlir::Location start = getLoc(s.getSourceRange().getEnd()); |
Also on the next line.
| LexicalScope LS{*this, start, builder.getInsertionBlock()}; | ||
| res = emitStmt(structuredBlock, /*useCurrentScope=*/true); | ||
|
|
||
| builder.create<Terminator>(end); |
There was a problem hiding this comment.
On first reading I didn't notice that Terminator was a template argument, and I was going to suggest mlir::acc::Terminator. Is the any convention for making this more obvious?
| const Stmt *structuredBlock) { | ||
| mlir::LogicalResult res = mlir::success(); | ||
|
|
||
| OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule()); |
| public: | ||
| OpenACCClauseCIREmitter(CIRGenModule &cgm) : cgm(cgm) {} | ||
|
|
||
| void VisitClauseList(llvm::ArrayRef<const OpenACCClause *> clauses) { |
There was a problem hiding this comment.
Is this necessary? It looks like it duplicates the OpenACCClauseVisitor implementation.
This LGTM. |
AllinLeeYL
pushed a commit
to AllinLeeYL/llvm-project
that referenced
this pull request
Apr 10, 2025
This patch adds some lowering code for Compute Constructs, plus the infrastructure to someday do clauses. Doing this requires adding the dialect to the CIRGenerator. This patch does not however implement/correctly initialize lowering from OpenACC-Dialect to anything lower however.
var-const
pushed a commit
to ldionne/llvm-project
that referenced
this pull request
Apr 17, 2025
This patch adds some lowering code for Compute Constructs, plus the infrastructure to someday do clauses. Doing this requires adding the dialect to the CIRGenerator. This patch does not however implement/correctly initialize lowering from OpenACC-Dialect to anything lower however.
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 patch adds some lowering code for Compute Constructs, plus the infrastructure to someday do clauses. Doing this requires adding the dialect to the CIRGenerator.
This patch does not however implement/correctly initialize lowering from OpenACC-Dialect to anything lower however.