[OpenACC][CIR] Initial patch to do OpenACC->IR lowering

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; }
 

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 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);

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

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);
 }

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
+}

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];
 

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
+}

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
+}