[OpenACC][CIR] Implement 'no_create' lowering for data

This lowering ends up being identical to 'create', except it is a
acc.nocreate for the start operation, and it doesn't permit modifier
list. This patch implements this by adding it to the list of permitted
handlers (along with compute), plus adds tests.

Author: erichkeane

clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp

@@ -939,17 +939,15 @@ class OpenACCClauseCIREmitter final
 
   void VisitNoCreateClause(const OpenACCNoCreateClause &clause) {
     if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
-                               mlir::acc::KernelsOp>) {
+                               mlir::acc::KernelsOp, mlir::acc::DataOp>) {
       for (const Expr *var : clause.getVarList())
         addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
             var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
             /*implicit=*/false);
     } else if constexpr (isCombinedType<OpTy>) {
       applyToComputeOp(clause);
     } else {
-      // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. data remains.
-      return clauseNotImplemented(clause);
+      llvm_unreachable("Unknown construct kind in VisitNoCreateClause");
     }
   }
 

clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c

@@ -187,4 +187,32 @@ void acc_data(int parmVar) {
   // CHECK-NEXT: } loc
   // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "parmVar"}
   // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
+
+#pragma acc data no_create(parmVar)
+  ;
+  // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+  // CHECK-NEXT: acc.data dataOperands(%[[NOCREATE1]] : !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "parmVar"}
+
+#pragma acc data no_create(parmVar) no_create(localVar1)
+  ;
+  // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+  // CHECK-NEXT: %[[NOCREATE2:.*]] = acc.nocreate varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
+  // CHECK-NEXT: acc.data dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "localVar1"}
+  // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "parmVar"}
+
+#pragma acc data no_create(parmVar, localVar1)
+  ;
+  // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+  // CHECK-NEXT: %[[NOCREATE2:.*]] = acc.nocreate varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
+  // CHECK-NEXT: acc.data dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "localVar1"}
+  // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "parmVar"}
 }