[clang] 2557f99 - [OpenACC][CIR] Implement 'no_create' lowering for data

[via cfe-commits]

Author: erichkeane
Date: 2025-06-27T09:27:59-07:00
New Revision: 2557f9946373148664ce90dc6da5a6677424ac2c

URL: https://github.com/llvm/llvm-project/commit/2557f9946373148664ce90dc6da5a6677424ac2c
DIFF: https://github.com/llvm/llvm-project/commit/2557f9946373148664ce90dc6da5a6677424ac2c.diff

LOG: [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.

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
    clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 1a078981a3a05..5652f03c92b13 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -939,7 +939,7 @@ 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,
@@ -947,9 +947,7 @@ class OpenACCClauseCIREmitter final
     } 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");
     }
   }
 

diff  --git a/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c b/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c
index 7a541a8c62ea6..068ca4a63e331 100644
--- a/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c
+++ b/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"}
 }