Merge commit 'e4d8e06f83f8' from llvm.org/main into next

Author: git apple-llvm automerger

clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp

@@ -922,7 +922,7 @@ class OpenACCClauseCIREmitter final
 
   void VisitDevicePtrClause(const OpenACCDevicePtrClause &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::DevicePtrOp>(
             var, mlir::acc::DataClause::acc_deviceptr, {},
@@ -932,7 +932,7 @@ class OpenACCClauseCIREmitter final
       applyToComputeOp(clause);
     } else {
       // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. data, declare remain.
+      // unreachable. declare remains.
       return clauseNotImplemented(clause);
     }
   }
@@ -953,7 +953,7 @@ class OpenACCClauseCIREmitter final
 
   void VisitPresentClause(const OpenACCPresentClause &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::PresentOp, mlir::acc::DeleteOp>(
             var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
@@ -962,14 +962,14 @@ class OpenACCClauseCIREmitter final
       applyToComputeOp(clause);
     } else {
       // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. data & declare remain.
+      // unreachable. declare remains.
       return clauseNotImplemented(clause);
     }
   }
 
   void VisitAttachClause(const OpenACCAttachClause &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::AttachOp, mlir::acc::DetachOp>(
             var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
@@ -978,7 +978,7 @@ class OpenACCClauseCIREmitter final
       applyToComputeOp(clause);
     } else {
       // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. data, enter data remain.
+      // unreachable. enter data remains.
       return clauseNotImplemented(clause);
     }
   }

clang/test/CIR/CodeGenOpenACC/data.c

@@ -3,6 +3,9 @@
 void acc_data(int cond) {
   // CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}) {
   // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+
+  int *ptr;
+  // CHECK-NEXT: %[[PTR:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptr"]
   // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
 
 #pragma acc data default(none)
@@ -221,5 +224,50 @@ void acc_data(int cond) {
   // CHECK-NEXT: acc.terminator
   // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}
 
+#pragma acc data deviceptr(ptr)
+  {}
+  // CHECK-NEXT: %[[DEV_PTR:.*]] = acc.deviceptr varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}
+  // CHECK-NEXT: acc.data dataOperands(%[[DEV_PTR]] : !cir.ptr<!cir.ptr<!s32i>>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+#pragma acc data deviceptr(ptr) device_type(radeon) async
+  {}
+  // CHECK-NEXT: %[[DEV_PTR:.*]] = acc.deviceptr varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}
+  // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[DEV_PTR]] : !cir.ptr<!cir.ptr<!s32i>>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+
+#pragma acc data present(cond)
+  {}
+  // CHECK-NEXT: %[[PRESENT:.*]] = acc.present varPtr(%[[COND]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "cond"}
+  // CHECK-NEXT: acc.data dataOperands(%[[PRESENT]] : !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.delete accPtr(%[[PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "cond"}
+
+#pragma acc data present(cond) device_type(radeon) async
+  {}
+  // CHECK-NEXT: %[[PRESENT:.*]] = acc.present varPtr(%[[COND]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) -> !cir.ptr<!s32i> {name = "cond"}
+  // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[PRESENT]] : !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.delete accPtr(%[[PRESENT]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_present>, name = "cond"}
+
+#pragma acc data attach(ptr)
+  {}
+  // CHECK-NEXT: %[[ATTACH:.*]] = acc.attach varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}
+  // CHECK-NEXT: acc.data dataOperands(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.detach accPtr(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_attach>, name = "ptr"}
+
+#pragma acc data attach(ptr) device_type(radeon) async
+  {}
+  // CHECK-NEXT: %[[ATTACH:.*]] = acc.attach varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}
+  // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.detach accPtr(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_attach>, name = "ptr"}
+
   // CHECK-NEXT: cir.return
 }