[OpenACC][CIR] Impl default/seq lowering for combined constructs

clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h

@@ -23,9 +23,25 @@ constexpr bool isOneOfTypes =
 template <typename ToTest, typename T>
 constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
 
+// Holds information for emitting clauses for a combined construct. We
+// instantiate the clause emitter with this type so that it can use
+// if-constexpr to specially handle these.
+template <typename CompOpTy> struct CombinedConstructClauseInfo {
+  using ComputeOpTy = CompOpTy;
+  ComputeOpTy computeOp;
+  mlir::acc::LoopOp loopOp;
+};
+
+template <typename ToTest> constexpr bool isCombinedType = false;
+template <typename T>
+constexpr bool isCombinedType<CombinedConstructClauseInfo<T>> = true;
+
 template <typename OpTy>
 class OpenACCClauseCIREmitter final
     : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> {
+  // Necessary for combined constructs.
+  template <typename FriendOpTy> friend class OpenACCClauseCIREmitter;
+
   OpTy &operation;
   CIRGen::CIRGenFunction &cgf;
   CIRGen::CIRGenBuilderTy &builder;
@@ -119,6 +135,26 @@ class OpenACCClauseCIREmitter final
     llvm_unreachable("unknown gang kind");
   }
 
+  template <typename U = void,
+            typename = std::enable_if_t<isCombinedType<OpTy>, U>>
+  void applyToLoopOp(const OpenACCClause &c) {
+    // TODO OpenACC: we have to set the insertion scope here correctly still.
+    OpenACCClauseCIREmitter<mlir::acc::LoopOp> loopEmitter{
+        operation.loopOp, cgf, builder, dirKind, dirLoc};
+    loopEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
+    loopEmitter.Visit(&c);
+  }
+
+  template <typename U = void,
+            typename = std::enable_if_t<isCombinedType<OpTy>, U>>
+  void applyToComputeOp(const OpenACCClause &c) {
+    // TODO OpenACC: we have to set the insertion scope here correctly still.
+    OpenACCClauseCIREmitter<typename OpTy::ComputeOpTy> computeEmitter{
+        operation.computeOp, cgf, builder, dirKind, dirLoc};
+    computeEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
+    computeEmitter.Visit(&c);
+  }
+
 public:
   OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
                           CIRGen::CIRGenBuilderTy &builder,
@@ -145,10 +181,10 @@ class OpenACCClauseCIREmitter final
       case OpenACCDefaultClauseKind::Invalid:
         break;
       }
+    } else if constexpr (isCombinedType<OpTy>) {
+      applyToComputeOp(clause);
     } else {
-      // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. Combined constructs remain.
-      return clauseNotImplemented(clause);
+      llvm_unreachable("Unknown construct kind in VisitDefaultClause");
     }
   }
 
@@ -175,9 +211,12 @@ class OpenACCClauseCIREmitter final
       // Nothing to do here, these constructs don't have any IR for these, as
       // they just modify the other clauses IR.  So setting of
       // `lastDeviceTypeValues` (done above) is all we need.
+    } else if constexpr (isCombinedType<OpTy>) {
+      // Nothing to do here either, combined constructs are just going to use
+      // 'lastDeviceTypeValues' to set the value for the child visitor.
     } else {
       // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. update, data, routine, combined constructs remain.
+      // unreachable. update, data, routine constructs remain.
       return clauseNotImplemented(clause);
     }
   }
@@ -334,9 +373,11 @@ class OpenACCClauseCIREmitter final
   void VisitSeqClause(const OpenACCSeqClause &clause) {
     if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
       operation.addSeq(builder.getContext(), lastDeviceTypeValues);
+    } else if constexpr (isCombinedType<OpTy>) {
+      applyToLoopOp(clause);
     } else {
       // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. Routine, Combined constructs remain.
+      // unreachable. Routine construct remains.
       return clauseNotImplemented(clause);
     }
   }

clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

@@ -109,6 +109,15 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpCombinedConstruct(
       builder.create<mlir::acc::YieldOp>(end);
     }
 
+    {
+      mlir::OpBuilder::InsertionGuard guardCase(builder);
+      CombinedConstructClauseInfo<Op> inf{computeOp, loopOp};
+      // We don't bother setting the insertion point, since the clause emitter
+      // is going to have to do this correctly.
+      makeClauseEmitter(inf, *this, builder, dirKind, dirLoc)
+          .VisitClauseList(clauses);
+    }
+
     builder.create<TermOp>(end);
   }
 

clang/test/CIR/CodeGenOpenACC/combined.cpp

@@ -22,13 +22,66 @@ extern "C" void acc_combined(int N) {
   // 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
+
+#pragma acc parallel loop default(none)
+  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: } attributes {defaultAttr = #acc<defaultvalue none>} loc
+
+#pragma acc serial loop default(present)
+  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: } attributes {defaultAttr = #acc<defaultvalue present>} loc
 
+#pragma acc kernels loop default(none)
+  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: } attributes {defaultAttr = #acc<defaultvalue none>} loc
+
+#pragma acc parallel loop seq
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.parallel combined(loop) {
+  // CHECK: acc.loop combined(parallel) {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop device_type(nvidia, radeon) seq
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.serial combined(loop) {
+  // CHECK: acc.loop combined(serial) {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc kernels loop seq device_type(nvidia, radeon)
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.kernels combined(loop) {
+  // CHECK: acc.loop combined(kernels) {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+  // CHECK: acc.terminator
   // CHECK-NEXT: } loc
+
 }

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

@@ -1,5 +1,4 @@
 // RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
-// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-llvm %s -o %t-cir.ll -verify
 
 void HelloWorld(int *A, int *B, int *C, int N) {
 
@@ -10,4 +9,11 @@ void HelloWorld(int *A, int *B, int *C, int N) {
 
 // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
 #pragma acc declare create(A)
+
+  // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: private}}
+#pragma acc parallel loop private(A)
+  for(int i = 0; i <5; ++i);
+  // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: async}}
+#pragma acc parallel loop async
+  for(int i = 0; i <5; ++i);
 }