diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index cc0f3b77c1a65..b7a73e2f62945 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -613,12 +613,39 @@ 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 +1122,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}
+}