From b843c0c8655fba8ce6d858d37c753bf0dbe304ee Mon Sep 17 00:00:00 2001 From: erichkeane Date: Tue, 8 Apr 2025 11:53:03 -0700 Subject: [PATCH 1/4] [OpenACC][CIR] Initial patch to do OpenACC->IR lowering 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. --- clang/include/clang/AST/OpenACCClause.h | 1 + clang/lib/CIR/CodeGen/CIRGenFunction.h | 10 +++ clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 72 ++++++++++++++++++- clang/lib/CIR/CodeGen/CIRGenerator.cpp | 2 + clang/test/CIR/CodeGenOpenACC/kernels.c | 30 ++++++++ .../openacc-not-implemented.cpp | 4 +- clang/test/CIR/CodeGenOpenACC/parallel.c | 29 ++++++++ clang/test/CIR/CodeGenOpenACC/serial.c | 30 ++++++++ 8 files changed, 174 insertions(+), 4 deletions(-) create mode 100644 clang/test/CIR/CodeGenOpenACC/kernels.c create mode 100644 clang/test/CIR/CodeGenOpenACC/parallel.c create mode 100644 clang/test/CIR/CodeGenOpenACC/serial.c 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 + mlir::LogicalResult + emitOpenACCComputeOp(mlir::Location start, + mlir::Location end, + llvm::ArrayRef 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 { + 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 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 +mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp( + mlir::Location start, mlir::Location end, + llvm::ArrayRef clauses, + const Stmt *structuredBlock) { + mlir::LogicalResult res = mlir::success(); + + OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule()); + + llvm::SmallVector retTy; + llvm::SmallVector operands; + auto op = builder.create(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(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( + start, end, s.clauses(), s.getStructuredBlock()); + case OpenACCDirectiveKind::Serial: + return emitOpenACCComputeOp( + start, end, s.clauses(), s.getStructuredBlock()); + case OpenACCDirectiveKind::Kernels: + return emitOpenACCComputeOp( + 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(); mlirContext->loadDialect(); + mlirContext->getOrLoadDialect(); cgm = std::make_unique( *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 +} From e59dd8f249489901144ef5ed6e3818e0584c9422 Mon Sep 17 00:00:00 2001 From: erichkeane Date: Tue, 8 Apr 2025 15:18:29 -0700 Subject: [PATCH 2/4] clang-format --- clang/lib/CIR/CodeGen/CIRGenFunction.h | 4 +-- clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 31 ++++++++++----------- 2 files changed, 17 insertions(+), 18 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 3628f0ea4510e..237c0acf85dbc 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -542,10 +542,10 @@ class CIRGenFunction : public CIRGenTypeCache { // other constructs, or at least be the base for construct emission. template mlir::LogicalResult - emitOpenACCComputeOp(mlir::Location start, - mlir::Location end, + emitOpenACCComputeOp(mlir::Location start, mlir::Location end, llvm::ArrayRef 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 cccf7ee826ea6..0b07e0b68039b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -25,28 +25,27 @@ using namespace mlir::acc; namespace { class OpenACCClauseCIREmitter final : public OpenACCClauseVisitor { - CIRGenModule &cgm; + CIRGenModule &cgm; - void clauseNotImplemented(const OpenACCClause &c) { - cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", - c.getClauseKind()); - } + void clauseNotImplemented(const OpenACCClause &c) { + cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind()); + } - public: - OpenACCClauseCIREmitter(CIRGenModule &cgm): cgm(cgm){} +public: + OpenACCClauseCIREmitter(CIRGenModule &cgm) : cgm(cgm) {} - void VisitClauseList(llvm::ArrayRef clauses) { - for (auto *clause : clauses) - Visit(clause); - } + void VisitClauseList(llvm::ArrayRef clauses) { + for (auto *clause : clauses) + Visit(clause); + } #define VISIT_CLAUSE(CN) \ - void Visit##CN##Clause(const OpenACC##CN##Clause &clause) { \ - clauseNotImplemented(clause); \ - } + void Visit##CN##Clause(const OpenACC##CN##Clause &clause) { \ + clauseNotImplemented(clause); \ + } #include "clang/Basic/OpenACCClauses.def" - }; -} +}; +} // namespace template mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp( From c8b6dd00436c59f0a0c3daf71598fccc9abd8ba5 Mon Sep 17 00:00:00 2001 From: erichkeane Date: Tue, 8 Apr 2025 16:34:52 -0700 Subject: [PATCH 3/4] Fixup feedback from Andy during code review --- clang/lib/CIR/CodeGen/CIRGenFunction.h | 2 +- clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 16 ++++++---------- 2 files changed, 7 insertions(+), 11 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 237c0acf85dbc..ce418c28609c1 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -540,7 +540,7 @@ class CIRGenFunction : public CIRGenTypeCache { // 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 + template mlir::LogicalResult emitOpenACCComputeOp(mlir::Location start, mlir::Location end, llvm::ArrayRef clauses, diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index 0b07e0b68039b..7a8879add784a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -34,11 +34,6 @@ class OpenACCClauseCIREmitter final public: OpenACCClauseCIREmitter(CIRGenModule &cgm) : cgm(cgm) {} - void VisitClauseList(llvm::ArrayRef clauses) { - for (auto *clause : clauses) - Visit(clause); - } - #define VISIT_CLAUSE(CN) \ void Visit##CN##Clause(const OpenACC##CN##Clause &clause) { \ clauseNotImplemented(clause); \ @@ -47,7 +42,7 @@ class OpenACCClauseCIREmitter final }; } // namespace -template +template mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp( mlir::Location start, mlir::Location end, llvm::ArrayRef clauses, @@ -55,6 +50,7 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp( mlir::LogicalResult res = mlir::success(); OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule()); + clauseEmitter.VisitClauseList(clauses); llvm::SmallVector retTy; llvm::SmallVector operands; @@ -64,17 +60,17 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp( mlir::OpBuilder::InsertionGuard guardCase(builder); builder.setInsertionPointToEnd(&block); - LexicalScope LS{*this, start, builder.getInsertionBlock()}; + LexicalScope ls{*this, start, builder.getInsertionBlock()}; res = emitStmt(structuredBlock, /*useCurrentScope=*/true); - builder.create(end); + builder.create(end); return res; } mlir::LogicalResult CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) { - auto start = getLoc(s.getSourceRange().getEnd()); - auto end = getLoc(s.getSourceRange().getEnd()); + mlir::Location start = getLoc(s.getSourceRange().getEnd()); + mlir::Location end = getLoc(s.getSourceRange().getEnd()); switch (s.getDirectiveKind()) { case OpenACCDirectiveKind::Parallel: From 9d5c270f93733901443a349dc32947f89763a33d Mon Sep 17 00:00:00 2001 From: erichkeane Date: Tue, 8 Apr 2025 16:36:32 -0700 Subject: [PATCH 4/4] fix terminator op check in serial.h --- clang/test/CIR/CodeGenOpenACC/serial.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c b/clang/test/CIR/CodeGenOpenACC/serial.c index 690fc7ad33e63..9897cd3d4e8d9 100644 --- a/clang/test/CIR/CodeGenOpenACC/serial.c +++ b/clang/test/CIR/CodeGenOpenACC/serial.c @@ -23,7 +23,7 @@ void acc_serial(void) { // CHECK-NEXT: } // cir.scope end: // CHECK-NEXT: } - // CHECK-NEXT:acc.terminator + // CHECK-NEXT:acc.yield // CHECK-NEXT:} // CHECK-NEXT: cir.return