[clang] [CIR] Represent Cleanups with Cleanup scopes (PR #180276)

Amr Hesham via cfe-commits cfe-commits at lists.llvm.org
Thu Feb 19 13:02:15 PST 2026


https://github.com/AmrDeveloper updated https://github.com/llvm/llvm-project/pull/180276

>From 54a0cd71cf19c8d5039a666c153f479242049f4c Mon Sep 17 00:00:00 2001
From: Amr Hesham <amr96 at programmer.net>
Date: Sun, 15 Feb 2026 20:07:01 +0100
Subject: [PATCH 1/4] [CIR] Represent EHCleanups with Cleanup scopes

---
 clang/lib/CIR/CodeGen/CIRGenCleanup.cpp       |  70 +++++++++-
 clang/lib/CIR/CodeGen/CIRGenCleanup.h         |  10 +-
 clang/lib/CIR/CodeGen/CIRGenFunction.cpp      |   6 +-
 clang/lib/CIR/CodeGen/CIRGenFunction.h        |  11 +-
 clang/test/CIR/CodeGen/array-dtor.cpp         |   2 +
 .../cleanup-scope-tmp-with-exception.cpp      |  31 +++++
 clang/test/CIR/CodeGen/cleanup-scope-tmp.cpp  |  31 +++++
 clang/test/CIR/CodeGen/destructors.cpp        |   2 +
 clang/test/CIR/CodeGen/dtors.cpp              |   2 +
 clang/test/CIR/CodeGen/nrvo.cpp               |   2 +
 clang/test/CIR/CodeGen/size-of-vla.cpp        |   2 +
 clang/test/CIR/CodeGen/try-catch-tmp.cpp      |   2 +
 clang/test/CIR/CodeGen/try-catch.cpp          |   2 +
 clang/test/CIR/CodeGen/vla.c                  |   2 +
 .../test/CIR/CodeGenOpenACC/declare-copy.cpp  | 123 ++++++++++-------
 .../CIR/CodeGenOpenACC/declare-copyin.cpp     | 101 ++++++++------
 .../CIR/CodeGenOpenACC/declare-copyout.cpp    | 128 +++++++++++-------
 .../CIR/CodeGenOpenACC/declare-create.cpp     | 126 ++++++++++-------
 .../CIR/CodeGenOpenACC/declare-deviceptr.cpp  |  54 ++++++--
 .../CodeGenOpenACC/declare-deviceresident.cpp | 127 ++++++++++-------
 .../test/CIR/CodeGenOpenACC/declare-link.cpp  |  28 +++-
 .../CIR/CodeGenOpenACC/declare-present.cpp    | 105 ++++++++------
 22 files changed, 654 insertions(+), 313 deletions(-)
 create mode 100644 clang/test/CIR/CodeGen/cleanup-scope-tmp-with-exception.cpp
 create mode 100644 clang/test/CIR/CodeGen/cleanup-scope-tmp.cpp

diff --git a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
index cc841c6474537..a53edaebf8f18 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
@@ -19,6 +19,7 @@
 #include "CIRGenCleanup.h"
 #include "CIRGenFunction.h"
 
+#include "clang/CIR/Dialect/IR/CIROpsEnums.h"
 #include "clang/CIR/MissingFeatures.h"
 
 using namespace clang;
@@ -144,8 +145,37 @@ void *EHScopeStack::pushCleanup(CleanupKind kind, size_t size) {
   bool isNormalCleanup = kind & NormalCleanup;
   bool isEHCleanup = kind & EHCleanup;
   bool isLifetimeMarker = kind & LifetimeMarker;
+  bool skipCleanupScope = false;
 
   assert(!cir::MissingFeatures::innermostEHScope());
+  cir::CleanupKind cleanupKind = cir::CleanupKind::All;
+  if (isEHCleanup && cgf->getLangOpts().Exceptions) {
+    cleanupKind =
+        isNormalCleanup ? cir::CleanupKind::All : cir::CleanupKind::EH;
+  } else {
+    if (isNormalCleanup)
+      cleanupKind = cir::CleanupKind::Normal;
+    else
+      skipCleanupScope = true;
+  }
+
+  cir::CleanupScopeOp cleanupScope = nullptr;
+  if (!skipCleanupScope) {
+    CIRGenBuilderTy &builder = cgf->getBuilder();
+    mlir::Location loc = builder.getUnknownLoc();
+    cleanupScope = cir::CleanupScopeOp::create(
+        builder, loc, cleanupKind,
+        /*bodyBuilder=*/
+        [&](mlir::OpBuilder &b, mlir::Location loc) {
+          // Terminations will be handled in popCleanup
+        },
+        /*cleanupBuilder=*/
+        [&](mlir::OpBuilder &b, mlir::Location loc) {
+          cir::YieldOp::create(b, loc);
+        });
+
+    builder.setInsertionPointToEnd(&cleanupScope.getBodyRegion().back());
+  }
 
   // Per C++ [except.terminate], it is implementation-defined whether none,
   // some, or all cleanups are called before std::terminate. Thus, when
@@ -157,7 +187,7 @@ void *EHScopeStack::pushCleanup(CleanupKind kind, size_t size) {
 
   EHCleanupScope *scope = new (buffer)
       EHCleanupScope(isNormalCleanup, isEHCleanup, size, branchFixups.size(),
-                     innermostNormalCleanup, innermostEHScope);
+                     cleanupScope, innermostNormalCleanup, innermostEHScope);
 
   if (isNormalCleanup)
     innermostNormalCleanup = stable_begin();
@@ -184,6 +214,18 @@ void EHScopeStack::popCleanup() {
   innermostNormalCleanup = cleanup.getEnclosingNormalCleanup();
   deallocate(cleanup.getAllocatedSize());
 
+  cir::CleanupScopeOp cleanupScope = cleanup.getCleanupScope();
+  if (cleanupScope) {
+    auto *block = &cleanupScope.getBodyRegion().back();
+    if (!block->mightHaveTerminator()) {
+      mlir::OpBuilder::InsertionGuard guard(cgf->getBuilder());
+      cgf->getBuilder().setInsertionPointToEnd(block);
+      cir::YieldOp::create(cgf->getBuilder(),
+                           cgf->getBuilder().getUnknownLoc());
+    }
+    cgf->getBuilder().setInsertionPointAfter(cleanupScope);
+  }
+
   // Destroy the cleanup.
   cleanup.destroy();
 
@@ -244,6 +286,8 @@ void CIRGenFunction::popCleanupBlock() {
   EHCleanupScope &scope = cast<EHCleanupScope>(*ehStack.begin());
   assert(scope.getFixupDepth() <= ehStack.getNumBranchFixups());
 
+  cir::CleanupScopeOp cleanupScope = scope.getCleanupScope();
+
   // Remember activation information.
   bool isActive = scope.isActive();
 
@@ -301,7 +345,19 @@ void CIRGenFunction::popCleanupBlock() {
     assert(!cir::MissingFeatures::ehCleanupScopeRequiresEHCleanup());
     ehStack.popCleanup();
     scope.markEmitted();
+
+    mlir::Block &block = cleanupScope.getCleanupRegion().back();
+    mlir::OpBuilder::InsertionGuard guard(builder);
+    builder.setInsertionPointToStart(&block);
+
     emitCleanup(*this, cleanup, cleanupFlags);
+
+    if (block.empty() ||
+        !block.back().hasTrait<mlir::OpTrait::IsTerminator>()) {
+      mlir::OpBuilder::InsertionGuard guardCase(builder);
+      builder.setInsertionPointToEnd(&block);
+      builder.createYield(cleanupScope.getLoc());
+    }
   } else {
     // Otherwise, the best approach is to thread everything through
     // the cleanup block and then try to clean up after ourselves.
@@ -363,8 +419,20 @@ void CIRGenFunction::popCleanupBlock() {
     ehStack.popCleanup();
     assert(ehStack.hasNormalCleanups() == hasEnclosingCleanups);
 
+    mlir::Block &block = cleanupScope.getCleanupRegion().back();
+
+    mlir::OpBuilder::InsertionGuard guard(builder);
+    builder.setInsertionPointToStart(&block);
+
     emitCleanup(*this, cleanup, cleanupFlags);
 
+    if (block.empty() ||
+        !block.back().hasTrait<mlir::OpTrait::IsTerminator>()) {
+      mlir::OpBuilder::InsertionGuard guardCase(builder);
+      builder.setInsertionPointToEnd(&block);
+      builder.createYield(cleanupScope.getLoc());
+    }
+
     // Append the prepared cleanup prologue from above.
     assert(!cir::MissingFeatures::cleanupAppendInsts());
 
diff --git a/clang/lib/CIR/CodeGen/CIRGenCleanup.h b/clang/lib/CIR/CodeGen/CIRGenCleanup.h
index a39eb02cf87d3..b345884e46f80 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCleanup.h
+++ b/clang/lib/CIR/CodeGen/CIRGenCleanup.h
@@ -19,6 +19,7 @@
 #include "EHScopeStack.h"
 #include "mlir/IR/Value.h"
 #include "clang/AST/StmtCXX.h"
+#include "clang/CIR/Dialect/IR/CIRDialect.h"
 
 namespace clang::CIRGen {
 
@@ -113,6 +114,8 @@ class alignas(EHScopeStack::ScopeStackAlignment) EHCleanupScope
   /// from this index onwards belong to this scope.
   unsigned fixupDepth = 0;
 
+  cir::CleanupScopeOp cleanupScope;
+
 public:
   /// Gets the size required for a lazy cleanup scope with the given
   /// cleanup-data requirements.
@@ -125,11 +128,12 @@ class alignas(EHScopeStack::ScopeStackAlignment) EHCleanupScope
   }
 
   EHCleanupScope(bool isNormal, bool isEH, unsigned cleanupSize,
-                 unsigned fixupDepth,
+                 unsigned fixupDepth, cir::CleanupScopeOp cleanupScope,
                  EHScopeStack::stable_iterator enclosingNormal,
                  EHScopeStack::stable_iterator enclosingEH)
       : EHScope(EHScope::Cleanup, enclosingEH),
-        enclosingNormal(enclosingNormal), fixupDepth(fixupDepth) {
+        enclosingNormal(enclosingNormal), fixupDepth(fixupDepth),
+        cleanupScope(cleanupScope) {
     cleanupBits.isNormalCleanup = isNormal;
     cleanupBits.isEHCleanup = isEH;
     cleanupBits.isActive = true;
@@ -168,6 +172,8 @@ class alignas(EHScopeStack::ScopeStackAlignment) EHCleanupScope
     return reinterpret_cast<EHScopeStack::Cleanup *>(getCleanupBuffer());
   }
 
+  cir::CleanupScopeOp getCleanupScope() { return cleanupScope; }
+
   static bool classof(const EHScope *scope) {
     return (scope->getKind() == Cleanup);
   }
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
index f8f674c4d2c21..72ba0856716c2 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
@@ -270,10 +270,12 @@ void CIRGenFunction::LexicalScope::cleanup() {
     // Leverage and defers to RunCleanupsScope's dtor and scope handling.
     applyCleanup();
 
+    mlir::Block *currentBlock = builder.getBlock();
+
     // If we now have one after `applyCleanup`, hook it up properly.
     if (!cleanupBlock && localScope->getCleanupBlock(builder)) {
       cleanupBlock = localScope->getCleanupBlock(builder);
-      cir::BrOp::create(builder, insPt->back().getLoc(), cleanupBlock);
+      cir::BrOp::create(builder, currentBlock->back().getLoc(), cleanupBlock);
       if (!cleanupBlock->mightHaveTerminator()) {
         mlir::OpBuilder::InsertionGuard guard(builder);
         builder.setInsertionPointToEnd(cleanupBlock);
@@ -309,7 +311,7 @@ void CIRGenFunction::LexicalScope::cleanup() {
     // End of any local scope != function
     // Ternary ops have to deal with matching arms for yielding types
     // and do return a value, it must do its own cir.yield insertion.
-    if (!localScope->isTernary() && !insPt->mightHaveTerminator()) {
+    if (!localScope->isTernary() && !currentBlock->mightHaveTerminator()) {
       !retVal ? cir::YieldOp::create(builder, localScope->endLoc)
               : cir::YieldOp::create(builder, localScope->endLoc, retVal);
     }
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index a21697ef8f1d4..cf1722e6cf6df 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -1054,13 +1054,10 @@ class CIRGenFunction : public CIRGenTypeCache {
     /// until this object is destroyed.
     void forceCleanup() {
       assert(performCleanup && "Already forced cleanup");
-      {
-        mlir::OpBuilder::InsertionGuard guard(cgf.getBuilder());
-        cgf.didCallStackSave = oldDidCallStackSave;
-        cgf.popCleanupBlocks(cleanupStackDepth);
-        performCleanup = false;
-        cgf.currentCleanupStackDepth = oldCleanupStackDepth;
-      }
+      cgf.didCallStackSave = oldDidCallStackSave;
+      cgf.popCleanupBlocks(cleanupStackDepth);
+      performCleanup = false;
+      cgf.currentCleanupStackDepth = oldCleanupStackDepth;
     }
   };
 
diff --git a/clang/test/CIR/CodeGen/array-dtor.cpp b/clang/test/CIR/CodeGen/array-dtor.cpp
index 4a3684efd0c20..9d6ef3a96024e 100644
--- a/clang/test/CIR/CodeGen/array-dtor.cpp
+++ b/clang/test/CIR/CodeGen/array-dtor.cpp
@@ -6,6 +6,8 @@
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
+// XFAIL: *
+
 struct S {
     ~S();
 };
diff --git a/clang/test/CIR/CodeGen/cleanup-scope-tmp-with-exception.cpp b/clang/test/CIR/CodeGen/cleanup-scope-tmp-with-exception.cpp
new file mode 100644
index 0000000000000..5ca5dd0ceeb9e
--- /dev/null
+++ b/clang/test/CIR/CodeGen/cleanup-scope-tmp-with-exception.cpp
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -std=c++20 -triple x86_64-unknown-linux-gnu -Wno-unused-value -fclangir -emit-cir %s -o %t.cir -fcxx-exceptions -fexceptions
+// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
+
+struct StructWithDestructor {
+  ~StructWithDestructor();
+  void procedure();
+};
+
+void cleanup_scope_with_without_body() { StructWithDestructor a; }
+
+// CIR: %[[A_ADDR:.*]] = cir.alloca !rec_StructWithDestructor, !cir.ptr<!rec_StructWithDestructor>, ["a"]
+// CIR: cir.cleanup.scope {
+// CIR:   cir.yield
+// CIR: } cleanup all {
+// CIR:   cir.call @_ZN20StructWithDestructorD1Ev(%[[A_ADDR]]) nothrow : (!cir.ptr<!rec_StructWithDestructor>) -> ()
+// CIR:   cir.yield
+// CIR: }
+
+void cleanup_scope_with_body_and_cleanup() {
+  StructWithDestructor a;
+  a.procedure();
+}
+
+// CIR: %[[A_ADDR:.*]] = cir.alloca !rec_StructWithDestructor, !cir.ptr<!rec_StructWithDestructor>, ["a"]
+// CIR: cir.cleanup.scope {
+// CIR:   cir.call @_ZN20StructWithDestructor9procedureEv(%[[A_ADDR]]) : (!cir.ptr<!rec_StructWithDestructor>) -> ()
+// CIR:   cir.yield
+// CIR: } cleanup all {
+// CIR:   cir.call @_ZN20StructWithDestructorD1Ev(%0) nothrow : (!cir.ptr<!rec_StructWithDestructor>) -> ()
+// CIR:   cir.yield
+// CIR: }
diff --git a/clang/test/CIR/CodeGen/cleanup-scope-tmp.cpp b/clang/test/CIR/CodeGen/cleanup-scope-tmp.cpp
new file mode 100644
index 0000000000000..a1fe64e206964
--- /dev/null
+++ b/clang/test/CIR/CodeGen/cleanup-scope-tmp.cpp
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -std=c++20 -triple x86_64-unknown-linux-gnu -Wno-unused-value -fclangir -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
+
+struct StructWithDestructor {
+  ~StructWithDestructor();
+  void procedure();
+};
+
+void cleanup_scope_with_without_body() { StructWithDestructor a; }
+
+// CIR: %[[A_ADDR:.*]] = cir.alloca !rec_StructWithDestructor, !cir.ptr<!rec_StructWithDestructor>, ["a"]
+// CIR: cir.cleanup.scope {
+// CIR:   cir.yield
+// CIR: } cleanup normal {
+// CIR:   cir.call @_ZN20StructWithDestructorD1Ev(%[[A_ADDR]]) nothrow : (!cir.ptr<!rec_StructWithDestructor>) -> ()
+// CIR:   cir.yield
+// CIR: }
+
+void cleanup_scope_with_body_and_cleanup() {
+  StructWithDestructor a;
+  a.procedure();
+}
+
+// CIR: %[[A_ADDR:.*]] = cir.alloca !rec_StructWithDestructor, !cir.ptr<!rec_StructWithDestructor>, ["a"]
+// CIR: cir.cleanup.scope {
+// CIR:   cir.call @_ZN20StructWithDestructor9procedureEv(%[[A_ADDR]]) : (!cir.ptr<!rec_StructWithDestructor>) -> ()
+// CIR:   cir.yield
+// CIR: } cleanup normal {
+// CIR:   cir.call @_ZN20StructWithDestructorD1Ev(%0) nothrow : (!cir.ptr<!rec_StructWithDestructor>) -> ()
+// CIR:   cir.yield
+// CIR: }
diff --git a/clang/test/CIR/CodeGen/destructors.cpp b/clang/test/CIR/CodeGen/destructors.cpp
index ec190f59b2f1d..7f439a51092a2 100644
--- a/clang/test/CIR/CodeGen/destructors.cpp
+++ b/clang/test/CIR/CodeGen/destructors.cpp
@@ -5,6 +5,8 @@
 // RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -mno-constructor-aliases -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
+// XFAIL: *
+
 void some_function() noexcept;
 
 struct out_of_line_destructor {
diff --git a/clang/test/CIR/CodeGen/dtors.cpp b/clang/test/CIR/CodeGen/dtors.cpp
index 895f800dac294..0bf72a83935eb 100644
--- a/clang/test/CIR/CodeGen/dtors.cpp
+++ b/clang/test/CIR/CodeGen/dtors.cpp
@@ -5,6 +5,8 @@
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -std=c++20 -mconstructor-aliases -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s --check-prefix=OGCG
 
+// XFAIL: *
+
 struct A {
   ~A();
 };
diff --git a/clang/test/CIR/CodeGen/nrvo.cpp b/clang/test/CIR/CodeGen/nrvo.cpp
index 0fc9b9ba54012..a4d4657b5be39 100644
--- a/clang/test/CIR/CodeGen/nrvo.cpp
+++ b/clang/test/CIR/CodeGen/nrvo.cpp
@@ -11,6 +11,8 @@
 // lowering isn't of interest for this test. We just need to see that the
 // copy constructor is elided without -fno-elide-constructors but not with it.
 
+// XFAIL: *
+
 struct S {
   S();
   int a;
diff --git a/clang/test/CIR/CodeGen/size-of-vla.cpp b/clang/test/CIR/CodeGen/size-of-vla.cpp
index bcaab27781aa3..5c42a38c90c49 100644
--- a/clang/test/CIR/CodeGen/size-of-vla.cpp
+++ b/clang/test/CIR/CodeGen/size-of-vla.cpp
@@ -5,6 +5,8 @@
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
+// XFAIL: *
+
 void vla_type_with_element_type_of_size_1() {
   unsigned long n = 10ul;
   unsigned long size = sizeof(bool[n]);
diff --git a/clang/test/CIR/CodeGen/try-catch-tmp.cpp b/clang/test/CIR/CodeGen/try-catch-tmp.cpp
index d43100c433e88..742db6e4e7327 100644
--- a/clang/test/CIR/CodeGen/try-catch-tmp.cpp
+++ b/clang/test/CIR/CodeGen/try-catch-tmp.cpp
@@ -3,6 +3,8 @@
 // RUN: %clang_cc1 -std=c++20 -triple x86_64-unknown-linux-gnu -fcxx-exceptions -fexceptions -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
+// XFAIL: *
+
 int division();
 
 void call_function_inside_try_catch_all() {
diff --git a/clang/test/CIR/CodeGen/try-catch.cpp b/clang/test/CIR/CodeGen/try-catch.cpp
index 81ac15eeb7016..0ba8e6a66a89e 100644
--- a/clang/test/CIR/CodeGen/try-catch.cpp
+++ b/clang/test/CIR/CodeGen/try-catch.cpp
@@ -5,6 +5,8 @@
 
 // TODO(cir): Reenable lowering to LLVM after new try-catch lowering is implemented.
 
+// XFAIL: *
+
 void empty_try_block_with_catch_all() {
   try {} catch (...) {}
 }
diff --git a/clang/test/CIR/CodeGen/vla.c b/clang/test/CIR/CodeGen/vla.c
index ce0cbee11add7..a70c721c2252e 100644
--- a/clang/test/CIR/CodeGen/vla.c
+++ b/clang/test/CIR/CodeGen/vla.c
@@ -5,6 +5,8 @@
 // RUN: %clang_cc1 -Wno-error=incompatible-pointer-types -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
+// XFAIL: *
+
 void f0(int len) {
   int arr[len];
 }
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp
index 1dd66826da96b..f593c2268c524 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp
@@ -53,13 +53,18 @@ struct Struct {
     // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
     //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "LocalInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHRCK-NEXT: }
   }
   void MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr);
 };
@@ -104,27 +109,38 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
 
 #pragma acc declare copy(alwaysout:LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.cleanup.scope {
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   } cleanup normal {
+    // CHECK-NEXT:     acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:     acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"}
+    // CHECK-NEXT:     acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
+    // CHECK-NEXT:     acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   }
+
+    // CHECK-NEXT:   cir.yield
+
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -160,36 +176,45 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "ArgHSEPtr[1:1]"}
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
     {
-      // CHECK-NEXT: cir.scope {
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.scope {
 #pragma acc declare copy(LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copy>, name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copy>, name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:     %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copy>, name = "LocalHSE"} 
+    // CHECK-NEXT:     %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "LocalInt"}
+    // CHECK-NEXT:     %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:     %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:     %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:     %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:     %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:     %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:     %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:     %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copy>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:     %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "LocalInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:     cir.cleanup.scope {
+    // CHECK-NEXT:       cir.call @do_thing
+    // CHECK-NEXT:       cir.yield
+    // CHECK-NEXT:     } cleanup normal {
+    // CHECK-NEXT:       acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:       acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, name = "LocalHSE"}
+    // CHECK-NEXT:       acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "LocalInt"}
+    // CHECK-NEXT:       acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:       cir.yield
+    // CHECK-NEXT:     }
     }
     // CHECK-NEXT: }
 
     // Make sure that cleanup gets put in the right scope.
     do_thing();
     // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- 
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "ArgInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT: cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
 }
 
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
index 9f8e78a58a97d..633c28ddf32c1 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
@@ -301,13 +301,18 @@ struct Struct {
     // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
     //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
   }
   void MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr);
 };
@@ -352,27 +357,38 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
 
 #pragma acc declare copyin(alwaysin:LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSEArr[1:1]"}
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.cleanup.scope {
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   } cleanup normal {
+    // CHECK-NEXT:     acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSE"}
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalInt"}
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:  }
+
+    // CHECK-NEXT:   cir.yield
+
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -408,7 +424,8 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
     {
-      // CHECK-NEXT: cir.scope {
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT: cir.scope {
 #pragma acc declare copyin(LocalHSE, LocalInt, LocalHSEArr[1:1])
     // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
     // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
@@ -423,21 +440,27 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_copyin>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_copyin>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
     }
     // CHECK-NEXT: }
 
     // Make sure that cleanup gets put in the right scope.
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- 
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
 }
 
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp
index 33e76a3b93e9c..d905de8301f76 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp
@@ -52,14 +52,19 @@ struct Struct {
     // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
     // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier always>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "LocalInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier always>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {modifiers = #acc<data_clause_modifier always>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
   }
   void MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr);
 };
@@ -104,27 +109,38 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
 
 #pragma acc declare copyout(alwaysout:LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "ArgInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
+ 
+    // CHECK-NEXT: cir.yield
+
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {modifiers = #acc<data_clause_modifier alwaysout>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -160,36 +176,46 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
     {
-      // CHECK-NEXT: cir.scope {
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT: cir.scope {
 #pragma acc declare copyout(LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copyout>, name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copyout>, name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copyout>, name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_copyout>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {name = "LocalHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {name = "LocalInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {name = "LocalHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) {name = "LocalInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
     }
     // CHECK-NEXT: }
 
     // Make sure that cleanup gets put in the right scope.
     do_thing();
     // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- 
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
-    // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT: cir.yield
+
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) {modifiers = #acc<data_clause_modifier always>, name = "ArgHSEPtr[1:1]"}
 }
 
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
index e5cf70190b849..7110447ac1f6b 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
@@ -301,14 +301,19 @@ struct Struct {
     // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
     // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {modifiers = #acc<data_clause_modifier zero>, name = "LocalHSEArr[1:1]"}
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSEPtr[1:1]"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSEArr[1:1]"}
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
   }
   void MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr);
 };
@@ -353,27 +358,38 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
 
 #pragma acc declare create(zero:LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {modifiers = #acc<data_clause_modifier zero>, name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {modifiers = #acc<data_clause_modifier zero>, name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {modifiers = #acc<data_clause_modifier zero>, name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {modifiers = #acc<data_clause_modifier zero>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "LocalHSEArr[1:1]"}
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT: cir.yield
+
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -409,36 +425,46 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {modifiers = #acc<data_clause_modifier zero>, name = "ArgHSEPtr[1:1]"}
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
     {
-      // CHECK-NEXT: cir.scope {
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT: cir.scope {
 #pragma acc declare create(LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_create>, name = "LocalHSEArr[1:1]"}
-    }
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_create>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
+  }
     // CHECK-NEXT: }
 
     // Make sure that cleanup gets put in the right scope.
     do_thing();
     // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT: cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
  
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp b/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp
index f6591f78aa225..f7559c8877b2e 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp
@@ -30,7 +30,12 @@ struct Struct {
     // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]], %[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
 
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]])
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
   }
   void MemFunc2(HasSideEffects *ArgHSE, int *ArgInt);
 };
@@ -59,12 +64,24 @@ void Struct::MemFunc2(HasSideEffects *ArgHSE, int *ArgInt) {
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
 
 #pragma acc declare deviceptr(LocalHSE, LocalInt)
-    // CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}
-    // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}
+    // CHECK-NEXT:   %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
     //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])
+    // CHECK-NEXT:   cir.cleanup.scope {
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   cleanup normal {
+    // CHECK-NEXT:     acc.declare_exit token(%[[ENTER2]])
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   }
+
+    // CHECK-NEXT:   cir.yield
+
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]])
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -84,14 +101,22 @@ void NormalFunc(HasSideEffects *ArgHSE, int *ArgInt) {
     // CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
     {
-      // CHECK-NEXT: cir.scope {
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT: cir.scope {
 #pragma acc declare deviceptr(LocalHSE, LocalInt)
-    // CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}
-    // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
+    // CHECK-NEXT:   %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}
+    // CHECK-NEXT:   %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
+
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]])
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 
     }
     // CHECK-NEXT: }
@@ -99,6 +124,9 @@ void NormalFunc(HasSideEffects *ArgHSE, int *ArgInt) {
     // Make sure that cleanup gets put in the right scope.
     do_thing();
     // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])
+    // CHECK-NEXT: cir.yield
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]])
+    // CHECK-NEXT:   cir.yield
 }
 
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp b/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp
index 6640c2581d24d..5d7c15647979b 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp
@@ -300,14 +300,19 @@ struct Struct {
     // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
     // CHECK-NEXT: %[[LOC_HSE_ARR_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSEPtr[1:1]"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSEArr[1:1]"}
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
   }
   void MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr);
 };
@@ -352,27 +357,38 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
 
 #pragma acc declare device_resident(LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   %[[LOC_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+
+    // CHECK-NEXT:   cir.cleanup.scope {
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   } cleanup normal {
+    // CHECK-NEXT:     acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSE"}
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalInt"}
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   }
 
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSEArr[1:1]"}
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:  cir.yield
+
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -408,36 +424,43 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ARG_HSE_PTR_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSEPtr[1:1]"}
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
     {
-      // CHECK-NEXT: cir.scope {
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT: cir.scope {
 #pragma acc declare device_resident(LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   %[[LOC_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.cleanup.scope {
+    // CHECK-NEXT:     cir.call @do_thing
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   } cleanup normal {
+    // CHECK-NEXT:     acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSE"}
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalInt"}
+    // CHECK-NEXT:     acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:     cir.yield
+    // CHECK-NEXT:   }
     }
     // CHECK-NEXT: }
 
     // Make sure that cleanup gets put in the right scope.
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- 
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_declare_device_resident>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
 }
-
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
index 5fc78167ce991..230d654fccdbe 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
@@ -143,8 +143,13 @@ struct Struct {
     // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
     //
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+    // 
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]])
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
   }
 
   void MemFunc2();
@@ -181,8 +186,13 @@ void Struct::MemFunc2() {
     // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr2[1:1]"}
     //
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+    // 
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]])
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -220,9 +230,13 @@ void NormalFunc() {
     // CHECK
 
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]])
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
     }
     // CHECK-NEXT: }
 
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-present.cpp b/clang/test/CIR/CodeGenOpenACC/declare-present.cpp
index 9c646d62a4f3c..85ceefac75cad 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-present.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-present.cpp
@@ -1,3 +1,4 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o %t.cir 
 // RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
 
 struct HasSideEffects {
@@ -52,14 +53,19 @@ struct Struct {
     // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
     // CHECK-NEXT: %[[LOC_HSE_ARR_PRESENT:.*]] = acc.present varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
     // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_present>, name = "ArgHSEPtr[1:1]"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_present>, name = "LocalHSEArr[1:1]"}
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup  normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_present>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_present>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
   }
   void MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr);
 };
@@ -104,27 +110,38 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
 
 #pragma acc declare present(LocalHSE, LocalInt, LocalHSEArr[1:1])
-    // CHECK-NEXT: %[[LOC_HSE_PRESENT:.*]] = acc.present varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
-    // CHECK-NEXT: %[[LOC_INT_PRESENT:.*]] = acc.present varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
-    // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
-    // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
-    // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
-    // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
-    // CHECK-NEXT: %[[LOC_HSE_ARR_PRESENT:.*]] = acc.present varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
-    // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   %[[LOC_HSE_PRESENT:.*]] = acc.present varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
+    // CHECK-NEXT:   %[[LOC_INT_PRESENT:.*]] = acc.present varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+    // CHECK-NEXT:   %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+    // CHECK-NEXT:   %[[IDX:.*]] = arith.constant 0 : i64
+    // CHECK-NEXT:   %[[STRIDE:.*]] = arith.constant 1 : i64
+    // CHECK-NEXT:   %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+    // CHECK-NEXT:   %[[LOC_HSE_ARR_PRESENT:.*]] = acc.present varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup  normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_present>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_present>, name = "LocalHSEArr[1:1]"}
-    //
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_present>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT: cir.yield
+
+    // CHECK-NEXT: cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_present>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
 }
 
 extern "C" void do_thing();
@@ -160,7 +177,8 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ARG_HSE_PTR_PRESENT:.*]] = acc.present varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSEPtr[1:1]"}
     // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
     {
-      // CHECK-NEXT: cir.scope {
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT: cir.scope {
 #pragma acc declare present(LocalHSE, LocalInt, LocalHSEArr[1:1])
     // CHECK-NEXT: %[[LOC_HSE_PRESENT:.*]] = acc.present varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} 
     // CHECK-NEXT: %[[LOC_INT_PRESENT:.*]] = acc.present varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
@@ -175,21 +193,28 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
     // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
 
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "LocalHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "LocalInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_present>, name = "LocalHSEArr[1:1]"}
-    }
+    // CHECK-NEXT: cir.cleanup.scope {
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_PRESENT]], %[[LOC_INT_PRESENT]], %[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "LocalHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "LocalInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[LOC_HSE_ARR_PRESENT]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) {dataClause = #acc<data_clause acc_present>, name = "LocalHSEArr[1:1]"}
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: }
+  }
     // CHECK-NEXT: }
 
     // Make sure that cleanup gets put in the right scope.
     do_thing();
-    // CHECK-NEXT: cir.call @do_thing
-    // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
+    // CHECK-NEXT:   cir.call @do_thing
+    // CHECK-NEXT:   cir.yield
+    // CHECK-NEXT: } cleanup normal {
+    // CHECK-NEXT:   acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_PRESENT]], %[[ARG_INT_PRESENT]], %[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
  
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "ArgHSE"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "ArgInt"}
-    // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_present>, name = "ArgHSEPtr[1:1]"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PRESENT]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_present>, name = "ArgHSE"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_INT_PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "ArgInt"}
+    // CHECK-NEXT:   acc.delete accPtr(%[[ARG_HSE_PTR_PRESENT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) {dataClause = #acc<data_clause acc_present>, name = "ArgHSEPtr[1:1]"}
 }
 

>From 831a3af700777eec3791e343403037b76558ff14 Mon Sep 17 00:00:00 2001
From: Amr Hesham <amr96 at programmer.net>
Date: Mon, 16 Feb 2026 19:11:26 +0100
Subject: [PATCH 2/4] Update Codegen tests

---
 clang/lib/CIR/CodeGen/CIRGenCleanup.cpp   |  2 +-
 clang/lib/CIR/CodeGen/CIRGenException.cpp | 10 ++++++
 clang/test/CIR/CodeGen/array-dtor.cpp     |  6 ++--
 clang/test/CIR/CodeGen/destructors.cpp    |  4 +--
 clang/test/CIR/CodeGen/size-of-vla.cpp    | 40 ++++++++++++++---------
 clang/test/CIR/CodeGen/try-catch-tmp.cpp  |  2 --
 clang/test/CIR/CodeGen/try-catch.cpp      |  2 --
 7 files changed, 38 insertions(+), 28 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
index a53edaebf8f18..7c957634c7425 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
@@ -171,7 +171,7 @@ void *EHScopeStack::pushCleanup(CleanupKind kind, size_t size) {
         },
         /*cleanupBuilder=*/
         [&](mlir::OpBuilder &b, mlir::Location loc) {
-          cir::YieldOp::create(b, loc);
+          // Terminations will be handled after emiting cleanup
         });
 
     builder.setInsertionPointToEnd(&cleanupScope.getBodyRegion().back());
diff --git a/clang/lib/CIR/CodeGen/CIRGenException.cpp b/clang/lib/CIR/CodeGen/CIRGenException.cpp
index 2d1d5e5e65f61..86aeb9f903172 100644
--- a/clang/lib/CIR/CodeGen/CIRGenException.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenException.cpp
@@ -12,6 +12,8 @@
 
 #include "CIRGenCXXABI.h"
 #include "CIRGenFunction.h"
+#include "mlir/IR/Block.h"
+#include "mlir/IR/Location.h"
 
 #include "clang/CIR/MissingFeatures.h"
 #include "llvm/Support/SaveAndRestore.h"
@@ -361,6 +363,7 @@ mlir::LogicalResult CIRGenFunction::emitCXXTryStmt(const CXXTryStmt &s) {
   for (unsigned i = 0; i != numHandlers; ++i) {
     const CXXCatchStmt *catchStmt = s.getHandler(i);
     mlir::Region *handler = &tryOp.getHandlerRegions()[i];
+    mlir::Location handlerLoc = getLoc(catchStmt->getCatchLoc());
 
     mlir::OpBuilder::InsertionGuard guard(builder);
     builder.setInsertionPointToStart(&handler->front());
@@ -391,6 +394,13 @@ mlir::LogicalResult CIRGenFunction::emitCXXTryStmt(const CXXTryStmt &s) {
 
     // Fall out through the catch cleanups.
     handlerScope.forceCleanup();
+
+    mlir::Block *block = &handler->getBlocks().back();
+    if (block->empty() ||
+        !block->back().hasTrait<mlir::OpTrait::IsTerminator>()) {
+      builder.setInsertionPointToEnd(block);
+      builder.createYield(handlerLoc);
+    }
   }
 
   return mlir::success();
diff --git a/clang/test/CIR/CodeGen/array-dtor.cpp b/clang/test/CIR/CodeGen/array-dtor.cpp
index 9d6ef3a96024e..0355a78729716 100644
--- a/clang/test/CIR/CodeGen/array-dtor.cpp
+++ b/clang/test/CIR/CodeGen/array-dtor.cpp
@@ -6,8 +6,6 @@
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
-// XFAIL: *
-
 struct S {
     ~S();
 };
@@ -47,10 +45,10 @@ void test_cleanup_array() {
 // CIR:   cir.return
 
 // LLVM: define{{.*}} void @_Z18test_cleanup_arrayv()
+// LLVM:   %[[ITER:.*]] = alloca ptr, i64 1
 // LLVM:   %[[ARRAY:.*]] = alloca [42 x %struct.S]
 // LLVM:   %[[START:.*]] = getelementptr %struct.S, ptr %[[ARRAY]], i32 0
 // LLVM:   %[[END:.*]] = getelementptr %struct.S, ptr %[[START]], i64 41
-// LLVM:   %[[ITER:.*]] = alloca ptr
 // LLVM:   store ptr %[[END]], ptr %[[ITER]]
 // LLVM:   br label %[[LOOP:.*]]
 // LLVM: [[COND:.*]]:
@@ -142,10 +140,10 @@ void multi_dimensional() {
 // CIR:       cir.return
 
 // LLVM:     define{{.*}} void @_Z17multi_dimensionalv()
+// LLVM:       %[[ITER:.*]] = alloca ptr, i64 1
 // LLVM:       %[[S:.*]] = alloca [3 x [5 x %struct.S]]
 // LLVM:       %[[START:.*]] = getelementptr %struct.S, ptr %[[S]], i32 0
 // LLVM:       %[[END:.*]] = getelementptr %struct.S, ptr %[[START]], i64 14
-// LLVM:       %[[ITER:.*]] = alloca ptr
 // LLVM:       store ptr %[[END]], ptr %[[ITER]]
 // LLVM:       br label %[[LOOP:.*]]
 // LLVM: [[COND:.*]]:
diff --git a/clang/test/CIR/CodeGen/destructors.cpp b/clang/test/CIR/CodeGen/destructors.cpp
index 7f439a51092a2..843a689ba8aca 100644
--- a/clang/test/CIR/CodeGen/destructors.cpp
+++ b/clang/test/CIR/CodeGen/destructors.cpp
@@ -5,8 +5,6 @@
 // RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -mno-constructor-aliases -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
-// XFAIL: *
-
 void some_function() noexcept;
 
 struct out_of_line_destructor {
@@ -100,6 +98,7 @@ void test_array_destructor() {
 // CIR:   }
 
 // LLVM: define{{.*}} void @_Z21test_array_destructorv()
+// LLVM:   %[[ARR_CUR:.*]] = alloca ptr, i64 1
 // LLVM:   %[[ARR:.*]] = alloca [5 x %struct.array_element]
 // LLVM:   %[[TMP:.*]] = alloca ptr
 // LLVM:   %[[ARR_PTR:.*]] = getelementptr %struct.array_element, ptr %[[ARR]], i32 0
@@ -118,7 +117,6 @@ void test_array_destructor() {
 // LLVM: [[INIT_LOOP_END]]:
 // LLVM:   %[[ARR_BEGIN:.*]] = getelementptr %struct.array_element, ptr %[[ARR]], i32 0
 // LLVM:   %[[ARR_END:.*]] = getelementptr %struct.array_element, ptr %[[ARR_BEGIN]], i64 4
-// LLVM:   %[[ARR_CUR:.*]] = alloca ptr
 // LLVM:   store ptr %[[ARR_END]], ptr %[[ARR_CUR]]
 // LLVM:   br label %[[DESTROY_LOOP_BODY:.*]]
 // LLVM: [[DESTROY_LOOP_NEXT:.*]]:
diff --git a/clang/test/CIR/CodeGen/size-of-vla.cpp b/clang/test/CIR/CodeGen/size-of-vla.cpp
index 5c42a38c90c49..268862285dc77 100644
--- a/clang/test/CIR/CodeGen/size-of-vla.cpp
+++ b/clang/test/CIR/CodeGen/size-of-vla.cpp
@@ -5,8 +5,6 @@
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
-// XFAIL: *
-
 void vla_type_with_element_type_of_size_1() {
   unsigned long n = 10ul;
   unsigned long size = sizeof(bool[n]);
@@ -67,25 +65,30 @@ void vla_expr_element_type_of_size_1() {
 
 // CIR: %[[N_ADDR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["n", init]
 // CIR: %[[SAVED_STACK_ADDR:.*]] = cir.alloca !cir.ptr<!u8i>, !cir.ptr<!cir.ptr<!u8i>>, ["saved_stack"]
+// CIR: %[[SIZE_ADDR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["size", init]
 // CIR: %[[CONST_10:.*]] = cir.const #cir.int<10> : !u64i
 // CIR: cir.store {{.*}} %[[CONST_10]], %[[N_ADDR]] : !u64i, !cir.ptr<!u64i>
 // CIR: %[[TMP_N:.*]] = cir.load {{.*}} %[[N_ADDR]] : !cir.ptr<!u64i>, !u64i
 // CIR: %[[STACK_SAVE:.*]] = cir.stacksave : !cir.ptr<!u8i>
 // CIR: cir.store {{.*}} %[[STACK_SAVE]], %[[SAVED_STACK_ADDR]] : !cir.ptr<!u8i>, !cir.ptr<!cir.ptr<!u8i>>
-// CIR: %[[ARR_ADDR:.*]] = cir.alloca !cir.bool, !cir.ptr<!cir.bool>, %[[TMP_N]] : !u64i, ["arr"]
-// CIR: %[[SIZE_ADDR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["size", init]
-// CIR: cir.store {{.*}} %[[TMP_N]], %[[SIZE_ADDR]] : !u64i, !cir.ptr<!u64i>
-// CIR: %[[TMP_SAVED_STACK:.*]] = cir.load {{.*}} %[[SAVED_STACK_ADDR]] : !cir.ptr<!cir.ptr<!u8i>>, !cir.ptr<!u8i>
-// CIR: cir.stackrestore %[[TMP_SAVED_STACK]] : !cir.ptr<!u8i>
+// CIR: cir.cleanup.scope {
+// CIR:   %[[ARR_ADDR:.*]] = cir.alloca !cir.bool, !cir.ptr<!cir.bool>, %[[TMP_N]] : !u64i, ["arr"]
+// CIR:   cir.store {{.*}} %[[TMP_N]], %[[SIZE_ADDR]] : !u64i, !cir.ptr<!u64i>
+// CIR:   cir.yield
+// CIR: } cleanup normal {
+// CIR:   %[[TMP_SAVED_STACK:.*]] = cir.load {{.*}} %[[SAVED_STACK_ADDR]] : !cir.ptr<!cir.ptr<!u8i>>, !cir.ptr<!u8i>
+// CIR:   cir.stackrestore %[[TMP_SAVED_STACK]] : !cir.ptr<!u8i>
+// CIR:   cir.yield
+// CIR: }
 
 // LLVM: %[[N_ADDR:.*]] = alloca i64, i64 1, align 8
 // LLVM: %[[SAVED_STACK_ADDR:.*]] = alloca ptr, i64 1, align 8
+// LLVM: %[[SIZE_ADDR:.*]] = alloca i64, i64 1, align 8
 // LLVM: store i64 10, ptr %[[N_ADDR]], align 8
 // LLVM: %[[TMP_N:.*]] = load i64, ptr %[[N_ADDR]], align 8
 // LLVM: %[[STACK_SAVE:.*]] = call ptr @llvm.stacksave.p0()
 // LLVM: store ptr %[[STACK_SAVE]], ptr %[[SAVED_STACK_ADDR]], align 8
 // LLVM: %[[ARR_ADDR:.*]] = alloca i8, i64 %[[TMP_N]], align 16
-// LLVM: %[[SIZE_ADDR:.*]] = alloca i64, i64 1, align 8
 // LLVM: store i64 %[[TMP_N]], ptr %[[SIZE_ADDR]], align 8
 // LLVM: %[[TMP_SAVED_STACK:.*]] = load ptr, ptr %[[SAVED_STACK_ADDR]], align 8
 // LLVM: call void @llvm.stackrestore.p0(ptr %[[TMP_SAVED_STACK]])
@@ -114,27 +117,32 @@ void vla_expr_element_type_int() {
 
 // CIR: %[[N_ADDR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["n", init]
 // CIR: %[[SAVED_STACK_ADDR:.*]] = cir.alloca !cir.ptr<!u8i>, !cir.ptr<!cir.ptr<!u8i>>, ["saved_stack"]
+// CIR: %[[SIZE_ADDR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["size", init]
 // CIR: %[[CONST_10:.*]] = cir.const #cir.int<10> : !u64i
 // CIR: cir.store {{.*}} %[[CONST_10]], %[[N_ADDR]] : !u64i, !cir.ptr<!u64i>
 // CIR: %[[TMP_N:.*]] = cir.load {{.*}} %[[N_ADDR]] : !cir.ptr<!u64i>, !u64i
 // CIR: %[[STACK_SAVE:.*]] = cir.stacksave : !cir.ptr<!u8i>
 // CIR: cir.store {{.*}} %[[STACK_SAVE]], %[[SAVED_STACK_ADDR]] : !cir.ptr<!u8i>, !cir.ptr<!cir.ptr<!u8i>>
-// CIR: %[[ARR_ADDR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, %[[TMP_N]] : !u64i, ["arr"]
-// CIR: %[[SIZE_ADDR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["size", init]
-// CIR: %[[CONST_4:.*]] = cir.const #cir.int<4> : !u64i
-// CIR: %[[SIZE:.*]] = cir.binop(mul, %[[CONST_4]], %[[TMP_N]]) nuw : !u64i
-// CIR: cir.store {{.*}} %[[SIZE]], %[[SIZE_ADDR]] : !u64i, !cir.ptr<!u64i>
-// CIR: %[[TMP_SAVED_STACK:.*]] = cir.load {{.*}} %[[SAVED_STACK_ADDR]] : !cir.ptr<!cir.ptr<!u8i>>, !cir.ptr<!u8i>
-// CIR: cir.stackrestore %[[TMP_SAVED_STACK]] : !cir.ptr<!u8i>
+// CIR: cir.cleanup.scope {
+// CIR:   %[[ARR_ADDR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, %[[TMP_N]] : !u64i, ["arr"]
+// CIR:   %[[CONST_4:.*]] = cir.const #cir.int<4> : !u64i
+// CIR:   %[[SIZE:.*]] = cir.binop(mul, %[[CONST_4]], %[[TMP_N]]) nuw : !u64i
+// CIR:   cir.store {{.*}} %[[SIZE]], %[[SIZE_ADDR]] : !u64i, !cir.ptr<!u64i>
+// CIR:   cir.yield
+// CIR: } cleanup normal {
+// CIR:   %[[TMP_SAVED_STACK:.*]] = cir.load {{.*}} %[[SAVED_STACK_ADDR]] : !cir.ptr<!cir.ptr<!u8i>>, !cir.ptr<!u8i>
+// CIR:   cir.stackrestore %[[TMP_SAVED_STACK]] : !cir.ptr<!u8i>
+// CIR:   cir.yield
+// CIR: }
 
 // LLVM: %[[N_ADDR:.*]] = alloca i64, i64 1, align 8
 // LLVM: %[[SAVED_STACK_ADDR:.*]] = alloca ptr, i64 1, align 8
+// LLVM: %[[SIZE_ADDR:.*]] = alloca i64, i64 1, align 8
 // LLVM: store i64 10, ptr %[[N_ADDR]], align 8
 // LLVM: %[[TMP_N:.*]] = load i64, ptr %[[N_ADDR]], align 8
 // LLVM: %[[STACK_SAVE:.*]] = call ptr @llvm.stacksave.p0()
 // LLVM: store ptr %[[STACK_SAVE]], ptr %[[SAVED_STACK_ADDR]], align 8
 // LLVM: %[[ARR_ADDR:.*]] = alloca i32, i64 %[[TMP_N]], align 16
-// LLVM: %[[SIZE_ADDR:.*]] = alloca i64, i64 1, align 8
 // LLVM: %[[SIZE:.*]] = mul nuw i64 4, %[[TMP_N]]
 // LLVM: store i64 %[[SIZE]], ptr %[[SIZE_ADDR]], align 8
 // LLVM: %[[TMP_SAVED_STACK:.*]] = load ptr, ptr %[[SAVED_STACK_ADDR]], align 8
diff --git a/clang/test/CIR/CodeGen/try-catch-tmp.cpp b/clang/test/CIR/CodeGen/try-catch-tmp.cpp
index 742db6e4e7327..d43100c433e88 100644
--- a/clang/test/CIR/CodeGen/try-catch-tmp.cpp
+++ b/clang/test/CIR/CodeGen/try-catch-tmp.cpp
@@ -3,8 +3,6 @@
 // RUN: %clang_cc1 -std=c++20 -triple x86_64-unknown-linux-gnu -fcxx-exceptions -fexceptions -emit-llvm %s -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
 
-// XFAIL: *
-
 int division();
 
 void call_function_inside_try_catch_all() {
diff --git a/clang/test/CIR/CodeGen/try-catch.cpp b/clang/test/CIR/CodeGen/try-catch.cpp
index 0ba8e6a66a89e..81ac15eeb7016 100644
--- a/clang/test/CIR/CodeGen/try-catch.cpp
+++ b/clang/test/CIR/CodeGen/try-catch.cpp
@@ -5,8 +5,6 @@
 
 // TODO(cir): Reenable lowering to LLVM after new try-catch lowering is implemented.
 
-// XFAIL: *
-
 void empty_try_block_with_catch_all() {
   try {} catch (...) {}
 }

>From 921f3069a67d80bc30d8f0ba4d7e1c6e2e93a526 Mon Sep 17 00:00:00 2001
From: Amr Hesham <amr96 at programmer.net>
Date: Wed, 18 Feb 2026 21:25:56 +0100
Subject: [PATCH 3/4] Address code review comments

---
 clang/lib/CIR/CodeGen/CIRGenCleanup.cpp   | 49 +++++++++--------------
 clang/lib/CIR/CodeGen/CIRGenCleanup.h     | 10 ++---
 clang/lib/CIR/CodeGen/CIRGenException.cpp |  1 +
 clang/test/CIR/CodeGen/array-dtor.cpp     |  2 +-
 clang/test/CIR/CodeGen/destructors.cpp    | 36 ++++++++++-------
 5 files changed, 47 insertions(+), 51 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
index 7c957634c7425..1451410e36f22 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
@@ -214,7 +214,7 @@ void EHScopeStack::popCleanup() {
   innermostNormalCleanup = cleanup.getEnclosingNormalCleanup();
   deallocate(cleanup.getAllocatedSize());
 
-  cir::CleanupScopeOp cleanupScope = cleanup.getCleanupScope();
+  cir::CleanupScopeOp cleanupScope = cleanup.getCleanupScopeOp();
   if (cleanupScope) {
     auto *block = &cleanupScope.getBodyRegion().back();
     if (!block->mightHaveTerminator()) {
@@ -256,13 +256,26 @@ bool EHScopeStack::requiresCatchOrCleanup() const {
   return false;
 }
 
-static void emitCleanup(CIRGenFunction &cgf, EHScopeStack::Cleanup *cleanup,
+static void emitCleanup(CIRGenFunction &cgf, cir::CleanupScopeOp cleanupScope,
+                        EHScopeStack::Cleanup *cleanup,
                         EHScopeStack::Cleanup::Flags flags) {
+  CIRGenBuilderTy &builder = cgf.getBuilder();
+  mlir::Block &block = cleanupScope.getCleanupRegion().back();
+
+  mlir::OpBuilder::InsertionGuard guard(builder);
+  builder.setInsertionPointToStart(&block);
+
   // Ask the cleanup to emit itself.
   assert(cgf.haveInsertPoint() && "expected insertion point");
   assert(!cir::MissingFeatures::ehCleanupActiveFlag());
   cleanup->emit(cgf, flags);
   assert(cgf.haveInsertPoint() && "cleanup ended with no insertion point?");
+
+  if (block.empty() || !block.back().hasTrait<mlir::OpTrait::IsTerminator>()) {
+    mlir::OpBuilder::InsertionGuard guardCase(builder);
+    builder.setInsertionPointToEnd(&block);
+    builder.createYield(cleanupScope.getLoc());
+  }
 }
 
 static mlir::Block *createNormalEntry(CIRGenFunction &cgf,
@@ -286,7 +299,8 @@ void CIRGenFunction::popCleanupBlock() {
   EHCleanupScope &scope = cast<EHCleanupScope>(*ehStack.begin());
   assert(scope.getFixupDepth() <= ehStack.getNumBranchFixups());
 
-  cir::CleanupScopeOp cleanupScope = scope.getCleanupScope();
+  cir::CleanupScopeOp cleanupScope = scope.getCleanupScopeOp();
+  assert(cleanupScope && "CleanupScopeOp is nullptr");
 
   // Remember activation information.
   bool isActive = scope.isActive();
@@ -345,19 +359,7 @@ void CIRGenFunction::popCleanupBlock() {
     assert(!cir::MissingFeatures::ehCleanupScopeRequiresEHCleanup());
     ehStack.popCleanup();
     scope.markEmitted();
-
-    mlir::Block &block = cleanupScope.getCleanupRegion().back();
-    mlir::OpBuilder::InsertionGuard guard(builder);
-    builder.setInsertionPointToStart(&block);
-
-    emitCleanup(*this, cleanup, cleanupFlags);
-
-    if (block.empty() ||
-        !block.back().hasTrait<mlir::OpTrait::IsTerminator>()) {
-      mlir::OpBuilder::InsertionGuard guardCase(builder);
-      builder.setInsertionPointToEnd(&block);
-      builder.createYield(cleanupScope.getLoc());
-    }
+    emitCleanup(*this, cleanupScope, cleanup, cleanupFlags);
   } else {
     // Otherwise, the best approach is to thread everything through
     // the cleanup block and then try to clean up after ourselves.
@@ -418,20 +420,7 @@ void CIRGenFunction::popCleanupBlock() {
     scope.markEmitted();
     ehStack.popCleanup();
     assert(ehStack.hasNormalCleanups() == hasEnclosingCleanups);
-
-    mlir::Block &block = cleanupScope.getCleanupRegion().back();
-
-    mlir::OpBuilder::InsertionGuard guard(builder);
-    builder.setInsertionPointToStart(&block);
-
-    emitCleanup(*this, cleanup, cleanupFlags);
-
-    if (block.empty() ||
-        !block.back().hasTrait<mlir::OpTrait::IsTerminator>()) {
-      mlir::OpBuilder::InsertionGuard guardCase(builder);
-      builder.setInsertionPointToEnd(&block);
-      builder.createYield(cleanupScope.getLoc());
-    }
+    emitCleanup(*this, cleanupScope, cleanup, cleanupFlags);
 
     // Append the prepared cleanup prologue from above.
     assert(!cir::MissingFeatures::cleanupAppendInsts());
diff --git a/clang/lib/CIR/CodeGen/CIRGenCleanup.h b/clang/lib/CIR/CodeGen/CIRGenCleanup.h
index b345884e46f80..aab678d032111 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCleanup.h
+++ b/clang/lib/CIR/CodeGen/CIRGenCleanup.h
@@ -19,7 +19,6 @@
 #include "EHScopeStack.h"
 #include "mlir/IR/Value.h"
 #include "clang/AST/StmtCXX.h"
-#include "clang/CIR/Dialect/IR/CIRDialect.h"
 
 namespace clang::CIRGen {
 
@@ -114,7 +113,8 @@ class alignas(EHScopeStack::ScopeStackAlignment) EHCleanupScope
   /// from this index onwards belong to this scope.
   unsigned fixupDepth = 0;
 
-  cir::CleanupScopeOp cleanupScope;
+  /// Cleanup scope op that represent the current scope in CIR
+  cir::CleanupScopeOp cleanupScopeOp;
 
 public:
   /// Gets the size required for a lazy cleanup scope with the given
@@ -128,12 +128,12 @@ class alignas(EHScopeStack::ScopeStackAlignment) EHCleanupScope
   }
 
   EHCleanupScope(bool isNormal, bool isEH, unsigned cleanupSize,
-                 unsigned fixupDepth, cir::CleanupScopeOp cleanupScope,
+                 unsigned fixupDepth, cir::CleanupScopeOp cleanupScopeOp,
                  EHScopeStack::stable_iterator enclosingNormal,
                  EHScopeStack::stable_iterator enclosingEH)
       : EHScope(EHScope::Cleanup, enclosingEH),
         enclosingNormal(enclosingNormal), fixupDepth(fixupDepth),
-        cleanupScope(cleanupScope) {
+        cleanupScopeOp(cleanupScopeOp) {
     cleanupBits.isNormalCleanup = isNormal;
     cleanupBits.isEHCleanup = isEH;
     cleanupBits.isActive = true;
@@ -172,7 +172,7 @@ class alignas(EHScopeStack::ScopeStackAlignment) EHCleanupScope
     return reinterpret_cast<EHScopeStack::Cleanup *>(getCleanupBuffer());
   }
 
-  cir::CleanupScopeOp getCleanupScope() { return cleanupScope; }
+  cir::CleanupScopeOp getCleanupScopeOp() { return cleanupScopeOp; }
 
   static bool classof(const EHScope *scope) {
     return (scope->getKind() == Cleanup);
diff --git a/clang/lib/CIR/CodeGen/CIRGenException.cpp b/clang/lib/CIR/CodeGen/CIRGenException.cpp
index 86aeb9f903172..59d87f40f2437 100644
--- a/clang/lib/CIR/CodeGen/CIRGenException.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenException.cpp
@@ -398,6 +398,7 @@ mlir::LogicalResult CIRGenFunction::emitCXXTryStmt(const CXXTryStmt &s) {
     mlir::Block *block = &handler->getBlocks().back();
     if (block->empty() ||
         !block->back().hasTrait<mlir::OpTrait::IsTerminator>()) {
+      mlir::OpBuilder::InsertionGuard guard(builder);
       builder.setInsertionPointToEnd(block);
       builder.createYield(handlerLoc);
     }
diff --git a/clang/test/CIR/CodeGen/array-dtor.cpp b/clang/test/CIR/CodeGen/array-dtor.cpp
index 0355a78729716..ba1e5f3acb202 100644
--- a/clang/test/CIR/CodeGen/array-dtor.cpp
+++ b/clang/test/CIR/CodeGen/array-dtor.cpp
@@ -140,7 +140,7 @@ void multi_dimensional() {
 // CIR:       cir.return
 
 // LLVM:     define{{.*}} void @_Z17multi_dimensionalv()
-// LLVM:       %[[ITER:.*]] = alloca ptr, i64 1
+// LLVM:       %[[ITER:.*]] = alloca ptr
 // LLVM:       %[[S:.*]] = alloca [3 x [5 x %struct.S]]
 // LLVM:       %[[START:.*]] = getelementptr %struct.S, ptr %[[S]], i32 0
 // LLVM:       %[[END:.*]] = getelementptr %struct.S, ptr %[[START]], i64 14
diff --git a/clang/test/CIR/CodeGen/destructors.cpp b/clang/test/CIR/CodeGen/destructors.cpp
index 843a689ba8aca..2a9c7336c5a39 100644
--- a/clang/test/CIR/CodeGen/destructors.cpp
+++ b/clang/test/CIR/CodeGen/destructors.cpp
@@ -79,24 +79,30 @@ void test_array_destructor() {
 // CIR:     %[[CMP:.*]] = cir.cmp(ne, %[[ARR_CUR]], %[[ARR_END]])
 // CIR:     cir.condition(%[[CMP]])
 // CIR:   }
-// CIR:   %[[FOUR:.*]] = cir.const #cir.int<4> : !u64i
-// CIR:   %[[BEGIN:.*]] = cir.cast array_to_ptrdecay %[[ARR]] : !cir.ptr<!cir.array<!rec_array_element x 5>>
-// CIR:   %[[END:.*]] = cir.ptr_stride %[[BEGIN]], %[[FOUR]] : (!cir.ptr<!rec_array_element>, !u64i)
-// CIR:   %[[ARR_PTR:.*]] = cir.alloca !cir.ptr<!rec_array_element>, !cir.ptr<!cir.ptr<!rec_array_element>>, ["__array_idx"]
-// CIR:   cir.store %[[END]], %[[ARR_PTR]]
-// CIR:   cir.do {
-// CIR:     %[[ARR_CUR:.*]] = cir.load{{.*}} %[[ARR_PTR]]
-// CIR:     cir.call @_ZN13array_elementD1Ev(%[[ARR_CUR]]) nothrow : (!cir.ptr<!rec_array_element>) -> ()
-// CIR:     %[[NEG_ONE:.*]] = cir.const #cir.int<-1> : !s64i
-// CIR:     %[[ARR_NEXT:.*]] = cir.ptr_stride %[[ARR_CUR]], %[[NEG_ONE]] : (!cir.ptr<!rec_array_element>, !s64i)
-// CIR:     cir.store %[[ARR_NEXT]], %[[ARR_PTR]]
+// CIR:   cir.cleanup.scope {
+// CIR:      cir.yield
+// CIR:   } cleanup normal {
+// CIR:     %[[FOUR:.*]] = cir.const #cir.int<4> : !u64i
+// CIR:     %[[BEGIN:.*]] = cir.cast array_to_ptrdecay %[[ARR]] : !cir.ptr<!cir.array<!rec_array_element x 5>>
+// CIR:     %[[END:.*]] = cir.ptr_stride %[[BEGIN]], %[[FOUR]] : (!cir.ptr<!rec_array_element>, !u64i)
+// CIR:     %[[ARR_PTR:.*]] = cir.alloca !cir.ptr<!rec_array_element>, !cir.ptr<!cir.ptr<!rec_array_element>>, ["__array_idx"]
+// CIR:     cir.store %[[END]], %[[ARR_PTR]]
+// CIR:     cir.do {
+// CIR:       %[[ARR_CUR:.*]] = cir.load{{.*}} %[[ARR_PTR]]
+// CIR:       cir.call @_ZN13array_elementD1Ev(%[[ARR_CUR]]) nothrow : (!cir.ptr<!rec_array_element>) -> ()
+// CIR:       %[[NEG_ONE:.*]] = cir.const #cir.int<-1> : !s64i
+// CIR:       %[[ARR_NEXT:.*]] = cir.ptr_stride %[[ARR_CUR]], %[[NEG_ONE]] : (!cir.ptr<!rec_array_element>, !s64i)
+// CIR:       cir.store %[[ARR_NEXT]], %[[ARR_PTR]]
+// CIR:       cir.yield
+// CIR:     } while {
+// CIR:       %[[ARR_CUR:.*]] = cir.load{{.*}} %[[ARR_PTR]]
+// CIR:       %[[CMP:.*]] = cir.cmp(ne, %[[ARR_CUR]], %[[BEGIN]])
+// CIR:       cir.condition(%[[CMP]])
+// CIR:     }
 // CIR:     cir.yield
-// CIR:   } while {
-// CIR:     %[[ARR_CUR:.*]] = cir.load{{.*}} %[[ARR_PTR]]
-// CIR:     %[[CMP:.*]] = cir.cmp(ne, %[[ARR_CUR]], %[[BEGIN]])
-// CIR:     cir.condition(%[[CMP]])
 // CIR:   }
 
+
 // LLVM: define{{.*}} void @_Z21test_array_destructorv()
 // LLVM:   %[[ARR_CUR:.*]] = alloca ptr, i64 1
 // LLVM:   %[[ARR:.*]] = alloca [5 x %struct.array_element]

>From cd3865cdbccf30d71da132cb5e918ce2878c15df Mon Sep 17 00:00:00 2001
From: Amr Hesham <amr96 at programmer.net>
Date: Thu, 19 Feb 2026 19:19:21 +0100
Subject: [PATCH 4/4] Address code review comments

---
 clang/lib/CIR/CodeGen/CIRGenCleanup.cpp  | 6 ++++--
 clang/lib/CIR/CodeGen/CIRGenFunction.cpp | 2 +-
 clang/test/CIR/CodeGen/array-dtor.cpp    | 2 +-
 clang/test/CIR/CodeGen/destructors.cpp   | 2 +-
 4 files changed, 7 insertions(+), 5 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
index 1451410e36f22..57f4f5f60a09d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCleanup.cpp
@@ -271,9 +271,11 @@ static void emitCleanup(CIRGenFunction &cgf, cir::CleanupScopeOp cleanupScope,
   cleanup->emit(cgf, flags);
   assert(cgf.haveInsertPoint() && "cleanup ended with no insertion point?");
 
-  if (block.empty() || !block.back().hasTrait<mlir::OpTrait::IsTerminator>()) {
+  mlir::Block &cleanupRegionLastBlock = cleanupScope.getCleanupRegion().back();
+  if (cleanupRegionLastBlock.empty() ||
+      !cleanupRegionLastBlock.back().hasTrait<mlir::OpTrait::IsTerminator>()) {
     mlir::OpBuilder::InsertionGuard guardCase(builder);
-    builder.setInsertionPointToEnd(&block);
+    builder.setInsertionPointToEnd(&cleanupRegionLastBlock);
     builder.createYield(cleanupScope.getLoc());
   }
 }
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
index 72ba0856716c2..52523dcccbe56 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
@@ -275,7 +275,7 @@ void CIRGenFunction::LexicalScope::cleanup() {
     // If we now have one after `applyCleanup`, hook it up properly.
     if (!cleanupBlock && localScope->getCleanupBlock(builder)) {
       cleanupBlock = localScope->getCleanupBlock(builder);
-      cir::BrOp::create(builder, currentBlock->back().getLoc(), cleanupBlock);
+      cir::BrOp::create(builder, insPt->back().getLoc(), cleanupBlock);
       if (!cleanupBlock->mightHaveTerminator()) {
         mlir::OpBuilder::InsertionGuard guard(builder);
         builder.setInsertionPointToEnd(cleanupBlock);
diff --git a/clang/test/CIR/CodeGen/array-dtor.cpp b/clang/test/CIR/CodeGen/array-dtor.cpp
index ba1e5f3acb202..fc05cf0b9f26c 100644
--- a/clang/test/CIR/CodeGen/array-dtor.cpp
+++ b/clang/test/CIR/CodeGen/array-dtor.cpp
@@ -45,7 +45,7 @@ void test_cleanup_array() {
 // CIR:   cir.return
 
 // LLVM: define{{.*}} void @_Z18test_cleanup_arrayv()
-// LLVM:   %[[ITER:.*]] = alloca ptr, i64 1
+// LLVM:   %[[ITER:.*]] = alloca ptr
 // LLVM:   %[[ARRAY:.*]] = alloca [42 x %struct.S]
 // LLVM:   %[[START:.*]] = getelementptr %struct.S, ptr %[[ARRAY]], i32 0
 // LLVM:   %[[END:.*]] = getelementptr %struct.S, ptr %[[START]], i64 41
diff --git a/clang/test/CIR/CodeGen/destructors.cpp b/clang/test/CIR/CodeGen/destructors.cpp
index 2a9c7336c5a39..0b17d1c7ab65b 100644
--- a/clang/test/CIR/CodeGen/destructors.cpp
+++ b/clang/test/CIR/CodeGen/destructors.cpp
@@ -104,7 +104,7 @@ void test_array_destructor() {
 
 
 // LLVM: define{{.*}} void @_Z21test_array_destructorv()
-// LLVM:   %[[ARR_CUR:.*]] = alloca ptr, i64 1
+// LLVM:   %[[ARR_CUR:.*]] = alloca ptr
 // LLVM:   %[[ARR:.*]] = alloca [5 x %struct.array_element]
 // LLVM:   %[[TMP:.*]] = alloca ptr
 // LLVM:   %[[ARR_PTR:.*]] = getelementptr %struct.array_element, ptr %[[ARR]], i32 0



More information about the cfe-commits mailing list