Author: Erich Keane Date: 2025-04-09T06:05:31-07:00 New Revision: 6e7c40b83de8290d081bcd72b1c57e30aeb28989
URL: https://github.com/llvm/llvm-project/commit/6e7c40b83de8290d081bcd72b1c57e30aeb28989 DIFF: https://github.com/llvm/llvm-project/commit/6e7c40b83de8290d081bcd72b1c57e30aeb28989.diff LOG: [OpenACC][CIR] Initial patch to do OpenACC->IR lowering (#134936) 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. Added: clang/test/CIR/CodeGenOpenACC/kernels.c clang/test/CIR/CodeGenOpenACC/parallel.c clang/test/CIR/CodeGenOpenACC/serial.c Modified: clang/include/clang/AST/OpenACCClause.h clang/lib/CIR/CodeGen/CIRGenFunction.h clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp clang/lib/CIR/CodeGen/CIRGenerator.cpp clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp Removed: ################################################################################ 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 9b4e364c40d9d..fb5ec6a868a1b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -570,6 +570,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 TermOp> + 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..7a8879add784a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -12,16 +12,79 @@ #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) {} + +#define VISIT_CLAUSE(CN) \ + void Visit##CN##Clause(const OpenACC##CN##Clause &clause) { \ + clauseNotImplemented(clause); \ + } +#include "clang/Basic/OpenACCClauses.def" +}; +} // namespace + +template <typename Op, typename TermOp> +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()); + clauseEmitter.VisitClauseList(clauses); + + 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<TermOp>(end); + return res; +} mlir::LogicalResult CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Compute Construct"); - return mlir::failure(); + mlir::Location start = getLoc(s.getSourceRange().getEnd()); + mlir::Location 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..9897cd3d4e8d9 --- /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.yield + // CHECK-NEXT:} + + // CHECK-NEXT: cir.return +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
