[OpenACC][CIR] Support for init and shutdown lowering

These two are very simple.  they don't require any clauses and don't
have an associated statement, so they have very simple output.  This
patch implements them, but none of the associated clauses.

Author: erichkeane

clang/lib/CIR/CodeGen/CIRGenFunction.h

@@ -571,6 +571,10 @@ class CIRGenFunction : public CIRGenTypeCache {
   //                         OpenACC Emission
   //===--------------------------------------------------------------------===//
 private:
+  template <typename Op>
+  mlir::LogicalResult
+  emitOpenACCOp(mlir::Location start,
+                llvm::ArrayRef<const OpenACCClause *> clauses);
   // Function to do the basic implementation of an operation with an Associated
   // Statement.  Models AssociatedStmtConstruct.
   template <typename Op, typename TermOp>

clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

@@ -94,6 +94,23 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
   return res;
 }
 
+template <typename Op>
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCOp(mlir::Location start,
+                              llvm::ArrayRef<const OpenACCClause *> clauses) {
+  mlir::LogicalResult res = mlir::success();
+
+  llvm::SmallVector<mlir::Type> retTy;
+  llvm::SmallVector<mlir::Value> operands;
+
+  // Clause-emitter must be here because it might modify operands.
+  OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
+  clauseEmitter.VisitClauseList(clauses);
+
+  builder.create<Op>(start, retTy, operands);
+  return res;
+}
+
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) {
   mlir::Location start = getLoc(s.getSourceRange().getEnd());
@@ -123,6 +140,17 @@ CIRGenFunction::emitOpenACCDataConstruct(const OpenACCDataConstruct &s) {
       start, end, s.clauses(), s.getStructuredBlock());
 }
 
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCInitConstruct(const OpenACCInitConstruct &s) {
+  mlir::Location start = getLoc(s.getSourceRange().getEnd());
+  return emitOpenACCOp<InitOp>(start, s.clauses());
+}
+mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct(
+    const OpenACCShutdownConstruct &s) {
+  mlir::Location start = getLoc(s.getSourceRange().getEnd());
+  return emitOpenACCOp<ShutdownOp>(start, s.clauses());
+}
+
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) {
   getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Loop Construct");
@@ -154,16 +182,6 @@ CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) {
   return mlir::failure();
 }
 mlir::LogicalResult
-CIRGenFunction::emitOpenACCInitConstruct(const OpenACCInitConstruct &s) {
-  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Init Construct");
-  return mlir::failure();
-}
-mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct(
-    const OpenACCShutdownConstruct &s) {
-  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Shutdown Construct");
-  return mlir::failure();
-}
-mlir::LogicalResult
 CIRGenFunction::emitOpenACCSetConstruct(const OpenACCSetConstruct &s) {
   getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Set Construct");
   return mlir::failure();

clang/test/CIR/CodeGenOpenACC/init.c

@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_init(void) {
+  // CHECK: cir.func @acc_init() {
+#pragma acc init
+// CHECK-NEXT: acc.init loc(#{{[a-zA-Z0-9]+}}){{$}}
+}

clang/test/CIR/CodeGenOpenACC/shutdown.c

@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_shutdown(void) {
+  // CHECK: cir.func @acc_shutdown() {
+#pragma acc shutdown
+// CHECK-NEXT: acc.shutdown loc(#{{[a-zA-Z0-9]+}}){{$}}
+}