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

This adds two clauses plus the infrastructure for emitting the clauses
on combined constructs.  Combined constructs require two operations, so
this makes sure we emit on the 'correct' one.  It DOES require that the
combined construct handling picks the correct one to put it on, AND sets
up the 'inserter' correctly, but these two clauses don't require an
inserter, so a future patch will get those.

Author: erichkeane

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