[OpenACC][CIR] Implement basic lowering for combined constructs

clang/lib/CIR/CodeGen/CIRGenFunction.h

@@ -718,6 +718,12 @@ class CIRGenFunction : public CIRGenTypeCache {
       SourceLocation dirLoc, llvm::ArrayRef<const OpenACCClause *> clauses,
       const Stmt *associatedStmt);
 
+  template <typename Op, typename TermOp>
+  mlir::LogicalResult emitOpenACCOpCombinedConstruct(
+      mlir::Location start, mlir::Location end, OpenACCDirectiveKind dirKind,
+      SourceLocation dirLoc, llvm::ArrayRef<const OpenACCClause *> clauses,
+      const Stmt *loopStmt);
+
 public:
   mlir::LogicalResult
   emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s);

clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

@@ -56,6 +56,65 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
   return res;
 }
 
+namespace {
+template <typename Op> struct CombinedType;
+template <> struct CombinedType<ParallelOp> {
+  static constexpr mlir::acc::CombinedConstructsType value =
+      mlir::acc::CombinedConstructsType::ParallelLoop;
+};
+template <> struct CombinedType<SerialOp> {
+  static constexpr mlir::acc::CombinedConstructsType value =
+      mlir::acc::CombinedConstructsType::SerialLoop;
+};
+template <> struct CombinedType<KernelsOp> {
+  static constexpr mlir::acc::CombinedConstructsType value =
+      mlir::acc::CombinedConstructsType::KernelsLoop;
+};
+} // namespace
+
+template <typename Op, typename TermOp>
+mlir::LogicalResult CIRGenFunction::emitOpenACCOpCombinedConstruct(
+    mlir::Location start, mlir::Location end, OpenACCDirectiveKind dirKind,
+    SourceLocation dirLoc, llvm::ArrayRef<const OpenACCClause *> clauses,
+    const Stmt *loopStmt) {
+  mlir::LogicalResult res = mlir::success();
+
+  llvm::SmallVector<mlir::Type> retTy;
+  llvm::SmallVector<mlir::Value> operands;
+
+  auto computeOp = builder.create<Op>(start, retTy, operands);
+  computeOp.setCombinedAttr(builder.getUnitAttr());
+  mlir::acc::LoopOp loopOp;
+
+  // First, emit the bodies of both operations, with the loop inside the body of
+  // the combined construct.
+  {
+    mlir::Block &block = computeOp.getRegion().emplaceBlock();
+    mlir::OpBuilder::InsertionGuard guardCase(builder);
+    builder.setInsertionPointToEnd(&block);
+
+    LexicalScope ls{*this, start, builder.getInsertionBlock()};
+    auto loopOp = builder.create<LoopOp>(start, retTy, operands);
+    loopOp.setCombinedAttr(mlir::acc::CombinedConstructsTypeAttr::get(
+        builder.getContext(), CombinedType<Op>::value));
+
+    {
+      mlir::Block &innerBlock = loopOp.getRegion().emplaceBlock();
+      mlir::OpBuilder::InsertionGuard guardCase(builder);
+      builder.setInsertionPointToEnd(&innerBlock);
+
+      LexicalScope ls{*this, start, builder.getInsertionBlock()};
+      res = emitStmt(loopStmt, /*useCurrentScope=*/true);
+
+      builder.create<mlir::acc::YieldOp>(end);
+    }
+
+    builder.create<TermOp>(end);
+  }
+
+  return res;
+}
+
 template <typename Op>
 Op CIRGenFunction::emitOpenACCOp(
     mlir::Location start, OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
@@ -170,8 +229,25 @@ CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) {
 
 mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
     const OpenACCCombinedConstruct &s) {
-  cgm.errorNYI(s.getSourceRange(), "OpenACC Combined Construct");
-  return mlir::failure();
+  mlir::Location start = getLoc(s.getSourceRange().getBegin());
+  mlir::Location end = getLoc(s.getSourceRange().getEnd());
+
+  switch (s.getDirectiveKind()) {
+  case OpenACCDirectiveKind::ParallelLoop:
+    return emitOpenACCOpCombinedConstruct<ParallelOp, mlir::acc::YieldOp>(
+        start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
+        s.getLoop());
+  case OpenACCDirectiveKind::SerialLoop:
+    return emitOpenACCOpCombinedConstruct<SerialOp, mlir::acc::YieldOp>(
+        start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
+        s.getLoop());
+  case OpenACCDirectiveKind::KernelsLoop:
+    return emitOpenACCOpCombinedConstruct<KernelsOp, mlir::acc::TerminatorOp>(
+        start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
+        s.getLoop());
+  default:
+    llvm_unreachable("invalid compute construct kind");
+  }
 }
 mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
     const OpenACCEnterDataConstruct &s) {

clang/test/CIR/CodeGenOpenACC/combined.cpp

@@ -0,0 +1,34 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+extern "C" void acc_combined(int N) {
+  // CHECK: cir.func @acc_combined(%[[ARG_N:.*]]: !s32i loc{{.*}}) {
+  // CHECK-NEXT: %[[ALLOCA_N:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["N", init]
+  // CHECK-NEXT: cir.store %[[ARG_N]], %[[ALLOCA_N]] : !s32i, !cir.ptr<!s32i>
+
+#pragma acc parallel loop
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.parallel combined(loop) {
+  // CHECK: acc.loop combined(parallel) {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc serial loop
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.serial combined(loop) {
+  // CHECK: acc.loop combined(serial) {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc kernels loop
+  for(unsigned I = 0; I < N; ++I);
+
+  // CHECK: acc.kernels combined(loop) {
+  // CHECK: acc.loop combined(kernels) {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK: acc.terminator
+  // CHECK-NEXT: } loc
+}

clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp

@@ -3,11 +3,10 @@
 
 void HelloWorld(int *A, int *B, int *C, int N) {
 
-// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Combined Construct}}
+// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Atomic Construct}}
 // expected-error@+1{{ClangIR code gen Not Yet Implemented: statement}}
-#pragma acc parallel loop
-  for (unsigned I = 0; I < N; ++I)
-    A[I] = B[I] + C[I];
+#pragma acc atomic
+  N = N + 1;
 
 // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
 #pragma acc declare create(A)