[clang] [OpenACC][CIR] 'update' construct lowering + a few clauses (PR #146378)

Mon, 30 Jun 2025 08:55:59 -0700 

https://github.com/erichkeane created 
https://github.com/llvm/llvm-project/pull/146378

The 'update' construct has 3 'var-list' clauses, device, self, and host. Each 
has a pretty simple data-operand type syntax in the IR, so this patch 
implements them as well.  At least one of those is required to be present on an 
'update', so we cannot do any lowering without them.

Note that 'self' and 'host' are aliases.

>From 87565b6e11a98d7fcab9262a95bb578c4465c6ad Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Mon, 30 Jun 2025 08:33:55 -0700
Subject: [PATCH] [OpenACC][CIR] 'update' construct lowering + a few clauses

The 'update' construct has 3 'var-list' clauses, device, self, and host.
Each has a pretty simple data-operand type syntax in the IR, so this
patch implements them as well.  At least one of those is required to be
present on an 'update', so we cannot do any lowering without them.

Note that 'self' and 'host' are aliases.
---
 clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 33 ++++++++-
 clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp   |  6 +-
 clang/test/CIR/CodeGenOpenACC/update.c        | 67 +++++++++++++++++++
 3 files changed, 101 insertions(+), 5 deletions(-)
 create mode 100644 clang/test/CIR/CodeGenOpenACC/update.c

diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp 
b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index cc0f3b77c1a65..6f9395df6928e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -613,12 +613,38 @@ class OpenACCClauseCIREmitter final
       } else {
         llvm_unreachable("var-list version of self shouldn't get here");
       }
+    } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      assert(!clause.isEmptySelfClause() && !clause.isConditionExprClause() &&
+             "var-list version of self required for update");
+      for (const Expr *var : clause.getVarList())
+        addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
+            var, mlir::acc::DataClause::acc_update_self, {},
+            /*structured=*/false, /*implicit=*/false);
     } else if constexpr (isCombinedType<OpTy>) {
       applyToComputeOp(clause);
     } else {
-      // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. update construct remains.
-      return clauseNotImplemented(clause);
+      llvm_unreachable("Unknown construct kind in VisitSelfClause");
+    }
+  }
+
+  void VisitHostClause(const OpenACCHostClause &clause){
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      for (const Expr *var : clause.getVarList())
+        addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
+            var, mlir::acc::DataClause::acc_update_host, {},
+            /*structured=*/false, /*implicit=*/false);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitHostClause");
+    }
+  }
+  void VisitDeviceClause(const OpenACCDeviceClause &clause){
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      for (const Expr *var : clause.getVarList())
+        addDataOperand<mlir::acc::UpdateDeviceOp>(
+            var, mlir::acc::DataClause::acc_update_device, {},
+            /*structured=*/false, /*implicit=*/false);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitDeviceClause");
     }
   }
 
@@ -1095,6 +1121,7 @@ EXPL_SPEC(mlir::acc::WaitOp)
 EXPL_SPEC(mlir::acc::HostDataOp)
 EXPL_SPEC(mlir::acc::EnterDataOp)
 EXPL_SPEC(mlir::acc::ExitDataOp)
+EXPL_SPEC(mlir::acc::UpdateOp)
 #undef EXPL_SPEC
 
 template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp 
b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index f3a635b7c83eb..5993056bf06ba 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -266,8 +266,10 @@ mlir::LogicalResult 
CIRGenFunction::emitOpenACCExitDataConstruct(
 
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
-  cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
-  return mlir::failure();
+  mlir::Location start = getLoc(s.getSourceRange().getBegin());
+  emitOpenACCOp<UpdateOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+                          s.clauses());
+  return mlir::success();
 }
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
diff --git a/clang/test/CIR/CodeGenOpenACC/update.c 
b/clang/test/CIR/CodeGenOpenACC/update.c
new file mode 100644
index 0000000000000..4e25a1df2a42b
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/update.c
@@ -0,0 +1,67 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir 
-fclangir %s -o - | FileCheck %s
+
+void acc_update(int parmVar, int *ptrParmVar) {
+  // CHECK: cir.func{{.*}} @acc_update(%[[ARG:.*]]: !s32i{{.*}}, 
%[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) {
+  // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", 
init]
+  // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, 
!cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]
+  // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>
+  // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, 
!cir.ptr<!cir.ptr<!s32i>>
+
+#pragma acc update device(parmVar)
+  // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]] : !cir.ptr<!s32i>)
+
+#pragma acc update device(parmVar, ptrParmVar)
+  // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : 
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", 
structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : 
!cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+
+#pragma acc update device(parmVar) device(ptrParmVar)
+  // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : 
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", 
structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : 
!cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+
+#pragma acc update host(parmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause 
acc_update_host>, name = "parmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to 
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc update host(parmVar, ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause 
acc_update_host>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : 
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = 
#acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : 
!cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) 
to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", 
structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to 
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc update host(parmVar) host(ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause 
acc_update_host>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : 
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = 
#acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : 
!cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) 
to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", 
structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to 
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to 
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar, ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : 
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = 
#acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : 
!cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) 
to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = 
#acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to 
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar) self(ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : 
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = 
#acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : 
!cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) 
to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = 
#acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to 
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar) device(ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : 
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", 
structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[UPD_DEV2]] : 
!cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to 
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause 
acc_update_self>, name = "parmVar", structured = false}
+}

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
  • [clang] [... Erich Keane via cfe-commits
    • [cla... via cfe-commits
    • [cla... Erich Keane via cfe-commits
    • [cla... Valentin Clement バレンタイン クレメン via cfe-commits
    • [cla... Erich Keane via cfe-commits
    • [cla... LLVM Continuous Integration via cfe-commits

Reply via email to